[clang] [llvm] [Clang][SYCL] Call sycl_kernel_launch with a kernel info type rather than a kernel name type. (PR #208958)
Tom Honermann via cfe-commits
cfe-commits at lists.llvm.org
Sat Jul 11 15:04:53 PDT 2026
https://github.com/tahonermann created https://github.com/llvm/llvm-project/pull/208958
Previously, implicit calls to the `sycl_kernel_launch()` function provided by the SYCL runtime library passed a kernel name type as the sole template argument and a name for the generated kernel entry point function as a string literal for the first function argument. This design had two issues:
- The generated kernel entry point function name was only communicated to the SYCL runtime library at run-time despite being known at compile-time.
- The interface left little room for future extension.
With this change, a kernel information type that contains both the kernel name type and the kernel entry point name as members is passed as the sole template argument and there is no longer a need to pass an implicit function argument. The kernel information type has the following structure where `kernel-info`, `kernel-name-type`, and `kernel-entry-point-name` are exposition only names.
```
struct kernel-info {
using kernel_name = kernel-name-type;
static constexpr const char kernel_entry_point_name[] = "kernel-entry-point-name";
};
```
The name of the kernel information type is formally unspecified and SYCL runtime authors should not write code that depends on it. Name mangling for these types is addressed by defining each such type as an explicit specialization of an implicitly declared class template named with the reserved identifier `__sycl_kernel_info`. The names generated for kernel entry point functions remain unspecified.
The kernel information type reserves room for future extension by allowing for additional compile-time information to be made available as additional members. For example, SYCL 2020 allows SYCL kernels to be defined with the `[[sycl::reqd_work_group_size(N)]]` attribute where `N` specifies a semantic constraint on how the SYCL runtime launches the kernel. Such constraints could be communicated by, for example, adding a `reqd_work_group_size` data member to the kernel information type without affecting the function call interface..
>From c072458a9113b573c14ec0608f1a85678e8116ac Mon Sep 17 00:00:00 2001
From: Tom Honermann <tom.honermann at intel.com>
Date: Sat, 11 Jul 2026 14:09:29 -0700
Subject: [PATCH] [Clang][SYCL] Call sycl_kernel_launch with a kernel info type
rather than a kernel name type.
Previously, implicit calls to the sycl_kernel_launch() function provided by
the SYCL runtime library passed a kernel name type as the sole template
argument and a name for the generated kernel entry point function as a
string literal for the first function argument. This design had two issues:
- The generated kernel entry point function name was only communicated to
the SYCL runtime library at run-time despite being known at compile-time.
- The interface left little room for future extension.
With this change, a kernel information type that contains both the kernel
name type and the kernel entry point name as members is passed as the sole
template argument and there is no longer a need to pass an implicit function
argument. The kernel information type has the following structure where
`kernel-info`, `kernel-name-type`, and `kernel-entry-point-name` are exposition
only names.
struct kernel-info {
using kernel_name = kernel-name-type;
static constexpr const char kernel_entry_point_name[] = "kernel-entry-point-name";
};
The name of the kernel information type is formally unspecified and SYCL
runtime authors should not write code that depends on it. Name mangling
for these types is addressed by defining each such type as an explicit
specialization of an implicitly declared class template named with the
reserved identifier `__sycl_kernel_info`. The names generated for kernel
entry point functions remain unspecified.
The kernel information type reserves room for future extension by allowing
for additional compile-time information to be made available as additional
members. For example, SYCL 2020 allows SYCL kernels to be defined with the
[[sycl::reqd_work_group_size(N)]] attribute where `N` specifies a semantic
constraint on how the SYCL runtime launches the kernel. Such constraints
could be communicated by, for example, adding a `reqd_work_group_size` data
member to the kernel information type without affecting the function call
interface..
---
clang/include/clang/AST/ASTNodeTraverser.h | 4 +-
clang/include/clang/AST/RecursiveASTVisitor.h | 1 +
clang/include/clang/AST/StmtSYCL.h | 24 +-
clang/include/clang/Basic/AttrDocs.td | 32 +-
clang/include/clang/Sema/ScopeInfo.h | 13 +-
clang/include/clang/Sema/SemaSYCL.h | 99 ++++--
clang/lib/Sema/SemaDecl.cpp | 58 ++--
clang/lib/Sema/SemaSYCL.cpp | 327 +++++++++++++++---
clang/lib/Sema/SemaTemplateInstantiate.cpp | 6 +-
clang/lib/Sema/TreeTransform.h | 13 +-
clang/lib/Serialization/ASTReaderStmt.cpp | 1 +
clang/lib/Serialization/ASTWriterStmt.cpp | 1 +
.../ast-dump-sycl-kernel-call-stmt.cpp | 100 +++---
.../ast-dump-sycl-kernel-entry-point.cpp | 6 +-
.../ASTSYCL/ast-print-sycl-kernel-call.cpp | 4 +-
clang/test/CodeGenSYCL/function-attrs.cpp | 2 +-
.../CodeGenSYCL/kernel-caller-entry-point.cpp | 57 ++-
.../sycl-kernel-entry-point-exceptions.cpp | 24 +-
clang/test/CodeGenSYCL/sycl-module-id.cpp | 4 +-
.../unique_stable_name_windows_diff.cpp | 4 +-
...-kernel-entry-point-attr-appertainment.cpp | 4 +-
...kernel-entry-point-attr-device-odr-use.cpp | 4 +-
.../sycl-kernel-entry-point-attr-grammar.cpp | 4 +-
...el-entry-point-attr-kernel-name-module.cpp | 4 +-
...ernel-entry-point-attr-kernel-name-pch.cpp | 4 +-
...cl-kernel-entry-point-attr-kernel-name.cpp | 4 +-
.../sycl-kernel-entry-point-attr-sfinae.cpp | 4 +-
.../sycl-kernel-entry-point-attr-this.cpp | 4 +-
.../SemaSYCL/sycl-kernel-launch-ms-compat.cpp | 12 +-
clang/test/SemaSYCL/sycl-kernel-launch.cpp | 204 +++++------
.../sycl-kernel-param-restrictions.cpp | 4 +-
libsycl/include/sycl/__impl/queue.hpp | 17 +-
32 files changed, 687 insertions(+), 362 deletions(-)
diff --git a/clang/include/clang/AST/ASTNodeTraverser.h b/clang/include/clang/AST/ASTNodeTraverser.h
index d184e03355077..061fd8d87199e 100644
--- a/clang/include/clang/AST/ASTNodeTraverser.h
+++ b/clang/include/clang/AST/ASTNodeTraverser.h
@@ -860,8 +860,10 @@ class ASTNodeTraverser
void
VisitUnresolvedSYCLKernelCallStmt(const UnresolvedSYCLKernelCallStmt *Node) {
Visit(Node->getOriginalStmt());
- if (Traversal != TK_IgnoreUnlessSpelledInSource)
+ if (Traversal != TK_IgnoreUnlessSpelledInSource) {
+ Visit(Node->getKernelInfoType());
Visit(Node->getKernelLaunchIdExpr());
+ }
}
void VisitOMPExecutableDirective(const OMPExecutableDirective *Node) {
diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h
index b77443f1fa0ab..cc62f4cd1fd17 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -3029,6 +3029,7 @@ DEF_TRAVERSE_STMT(SYCLUniqueStableNameExpr, {
DEF_TRAVERSE_STMT(UnresolvedSYCLKernelCallStmt, {
if (getDerived().shouldVisitImplicitCode()) {
TRY_TO(TraverseStmt(S->getOriginalStmt()));
+ TRY_TO(TraverseType(S->getKernelInfoType()));
TRY_TO(TraverseStmt(S->getKernelLaunchIdExpr()));
ShouldVisitChildren = false;
}
diff --git a/clang/include/clang/AST/StmtSYCL.h b/clang/include/clang/AST/StmtSYCL.h
index 79ac88532e143..75acf1bb9c7f6 100644
--- a/clang/include/clang/AST/StmtSYCL.h
+++ b/clang/include/clang/AST/StmtSYCL.h
@@ -103,27 +103,39 @@ class UnresolvedSYCLKernelCallStmt : public Stmt {
private:
Stmt *OriginalStmt = nullptr;
+
+ // The kernel information type to use in synthesized calls to the SYCL
+ // runtime library. This type is a specialization of the class template
+ // returned by SemaSYCL::GetSYCLKernelInfoClassTemplate() specialized for
+ // the kernel name type. This will be a dependent type if the kernel name
+ // type is dependent.
+ QualType KernelInfoType;
+
// 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)
+ UnresolvedSYCLKernelCallStmt(CompoundStmt *CS, QualType InfoType,
+ Expr *IdExpr)
: Stmt(UnresolvedSYCLKernelCallStmtClass), OriginalStmt(CS),
- KernelLaunchIdExpr(IdExpr) {}
+ KernelInfoType(InfoType), KernelLaunchIdExpr(IdExpr) {}
void setOriginalStmt(CompoundStmt *CS) { OriginalStmt = CS; }
+ void setKernelInfoType(QualType InfoType) { KernelInfoType = InfoType; }
+
void setKernelLaunchIdExpr(Expr *IdExpr) { KernelLaunchIdExpr = IdExpr; }
public:
static UnresolvedSYCLKernelCallStmt *Create(const ASTContext &C,
- CompoundStmt *CS, Expr *IdExpr) {
- return new (C) UnresolvedSYCLKernelCallStmt(CS, IdExpr);
+ CompoundStmt *CS,
+ QualType InfoType, Expr *IdExpr) {
+ return new (C) UnresolvedSYCLKernelCallStmt(CS, InfoType, IdExpr);
}
static UnresolvedSYCLKernelCallStmt *CreateEmpty(const ASTContext &C) {
- return new (C) UnresolvedSYCLKernelCallStmt(nullptr, nullptr);
+ return new (C) UnresolvedSYCLKernelCallStmt(nullptr, QualType(), nullptr);
}
CompoundStmt *getOriginalStmt() { return cast<CompoundStmt>(OriginalStmt); }
@@ -131,6 +143,8 @@ class UnresolvedSYCLKernelCallStmt : public Stmt {
return cast<CompoundStmt>(OriginalStmt);
}
+ QualType getKernelInfoType() const { return KernelInfoType; }
+
Expr *getKernelLaunchIdExpr() { return KernelLaunchIdExpr; }
const Expr *getKernelLaunchIdExpr() const { return KernelLaunchIdExpr; }
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 05e4cb0870652..1a2d9244e064d 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -701,13 +701,23 @@ follows.
namespace sycl {
class handler {
- template<typename KernelName, typename... Ts>
- void sycl_kernel_launch(const char* kernelSymbol, Ts&&... kernelArgs) {
+ template<typename KernelInfo, typename... Ts>
+ void sycl_kernel_launch(Ts&&... kernelArgs) {
// This code will run on the host and is responsible for calling functions
// appropriate for the desired offload backend (OpenCL, CUDA, HIP,
// Level Zero, etc...) to copy the kernel arguments denoted by kernelArgs
// to a device and to schedule an invocation of the offload kernel entry
- // point denoted by kernelSymbol with the copied arguments.
+ // point with the copied arguments.
+ //
+ // The KernelInfo template parameter denotes an unspecified type which
+ // contains metadata for the kernel. It has the following members:
+ // - kernel_name: a type alias for the kernel name type. This corresponds
+ // to the KernelName template parameter of SYCL kernel invocation
+ // functions.
+ // - kernel_entry_point_name: a static constexpr array of char containing
+ // the name of the generated offload kernel entry point function. The
+ // array is suitable for use in constant expressions, including non-type
+ // template arguments.
}
template<typename KernelName, typename KernelType>
@@ -823,12 +833,16 @@ There are a few items worthy of note:
copyability and kernel parameter requirements.
Within ``single_task()``, the call to ``kernel_entry_point()`` is effectively
-replaced with a synthesized call to a ''sycl_kernel_launch`` template that
+replaced with a synthesized call to a ``sycl_kernel_launch`` template that
looks approximately as follows.
.. code-block:: c++
- sycl_kernel_launch<KN>("sycl-kernel-caller-for-KN", kernelFunc);
+ struct kernel-info {
+ using kernel_name = KN;
+ static constexpr const char kernel_entry_point_name[] = "sycl-kernel-caller-for-KN";
+ };
+ sycl_kernel_launch<KN>(kernelFunc);
There are a few items worthy of note:
@@ -844,10 +858,10 @@ There are a few items worthy of note:
#. The ``sycl_kernel_launch`` template is expected to be provided by the SYCL
library implementation. It is responsible for copying the kernel arguments
to device memory and for scheduling execution of the generated offload
- kernel entry point identified by the symbol name passed as the first
- function argument. ``sycl-kernel-caller-for-KN`` is substituted above for
- the actual symbol name that would be generated for the offload kernel entry
- point.
+ kernel entry point identified by the ``kernel_entry_point_name`` member of
+ the first template parameter. ``sycl-kernel-caller-for-KN`` is substituted
+ above for the actual symbol name that would be generated for the offload
+ kernel entry point.
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
diff --git a/clang/include/clang/Sema/ScopeInfo.h b/clang/include/clang/Sema/ScopeInfo.h
index 7e4d3f2e0d1cb..fb92ee0d2ef49 100644
--- a/clang/include/clang/Sema/ScopeInfo.h
+++ b/clang/include/clang/Sema/ScopeInfo.h
@@ -23,6 +23,7 @@
#include "clang/Basic/SourceLocation.h"
#include "clang/Sema/CleanupInfo.h"
#include "clang/Sema/DeclSpec.h"
+#include "clang/Sema/SemaSYCL.h"
#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/DenseMapInfo.h"
#include "llvm/ADT/MapVector.h"
@@ -245,9 +246,15 @@ 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;
+ // Implicit calls to the SYCL runtime library are synthesized for functions
+ // declared with the sycl_kernel_entry_point attribute. These calls require
+ // types and lookup results for SYCL runtime library declared functions that
+ // are produced at the start of a function definition. SYCLKernelASTFragments
+ // caches these artifacts for later use to construct a SYCLKernelCallStmt or
+ // UnresolvedSYCLKernelCallStmt. If production of these artifacts fails, a
+ // diagnostic will have been issued and a disengaged optional stored.
+ std::optional<SemaSYCL::SYCLKernelCallStmtASTFragments>
+ SYCLKernelASTFragments;
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 4980aa44c3012..fff3c65983365 100644
--- a/clang/include/clang/Sema/SemaSYCL.h
+++ b/clang/include/clang/Sema/SemaSYCL.h
@@ -18,6 +18,7 @@
#include "clang/Basic/SourceLocation.h"
#include "clang/Sema/Ownership.h"
#include "clang/Sema/SemaBase.h"
+#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/DenseSet.h"
namespace clang {
@@ -25,6 +26,16 @@ class Decl;
class ParsedAttr;
class SemaSYCL : public SemaBase {
+ /// SYCLKernelInfoClassTemplate caches the class template declaration built
+ /// by GetSYCLKernelInfoClassTemplate().
+ ClassTemplateDecl *SYCLKernelInfoClassTemplate = nullptr;
+
+ /// SYCLKernelInfoClassTemplateSpecializations maps SYCL kernel name types
+ /// to their corresponding synthesized explicit specialization declarations
+ /// of the SYCL kernel info class template.
+ llvm::DenseMap<CanQualType, ClassTemplateSpecializationDecl *>
+ SYCLKernelInfoClassTemplateSpecializations;
+
public:
SemaSYCL(Sema &S);
@@ -71,31 +82,69 @@ class SemaSYCL : public SemaBase {
void CheckSYCLExternalFunctionDecl(FunctionDecl *FD);
void CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD);
- /// Builds an expression for the lookup of a 'sycl_kernel_launch' template
- /// with 'KernelName' as an explicit template argument. Lookup is performed
- /// as if from the first statement of the body of 'FD' and thus requires
- /// searching the scopes that exist at parse time. This function therefore
- /// requires the current semantic context to be the definition of 'FD'. In a
- /// dependent context, the returned expression will be an UnresolvedLookupExpr
- /// or an UnresolvedMemberExpr. In a non-dependent context, the returned
- /// expression will be a DeclRefExpr or MemberExpr. If lookup fails, a null
- /// error result is returned. The resulting expression is intended to be
- /// passed as the 'LaunchIdExpr' argument in a call to either
- /// BuildSYCLKernelCallStmt() or BuildUnresolvedSYCLKernelCallStmt() after
- /// the function body has been parsed.
- ExprResult BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD, QualType KernelName);
-
- /// Builds a SYCLKernelCallStmt to wrap 'Body' and to be used as the body of
- /// 'FD'. 'LaunchIdExpr' specifies the lookup result returned by a previous
- /// call to BuildSYCLKernelLaunchIdExpr().
- StmtResult BuildSYCLKernelCallStmt(FunctionDecl *FD, CompoundStmt *Body,
- Expr *LaunchIdExpr);
-
- /// Builds an UnresolvedSYCLKernelCallStmt to wrap 'Body'. 'LaunchIdExpr'
- /// specifies the lookup result returned by a previous call to
- /// BuildSYCLKernelLaunchIdExpr().
- StmtResult BuildUnresolvedSYCLKernelCallStmt(CompoundStmt *Body,
- Expr *LaunchIdExpr);
+ /// GetSYCLKernelInfoClassTemplate builds a class template used to pass SYCL
+ /// kernel information in synthesized calls to SYCL runtime library functions.
+ /// The class template has an unspecified name, is an incomplete class
+ /// (explicit specializations are synthesized as needed), built on first call,
+ /// and cached for return in subsequent calls.
+ ClassTemplateDecl *GetSYCLKernelInfoClassTemplate();
+
+ /// GetSYCLKernelInfoClassSpecializationType returns a type for the class
+ /// template returned by GetSYCLKernelInfoClassTemplate() specialized for
+ /// 'KNT'. 'KNT' may be a dependent type. If 'KNT' is not a dependent type,
+ /// an explicit specialization declaration for the specialization is
+ /// synthesized and a canonical type returned. A definition for the explicit
+ /// specialization is synthesized when SYCLKernelCallStmt is constructed for
+ /// 'KNT'.
+ QualType GetSYCLKernelInfoClassSpecializationType(QualType KNT);
+
+ /// SYCLKernelCallStmtASTFragments holds portions of the AST used for lookup
+ /// of SYCL runtime library functions. These are constructed at the beginning
+ /// of a function definition and later used in synthesized calls to the
+ /// SYCL runtime library. In a dependent context, these are used to construct
+ /// an UnresolvedSYCLKernelCallStmt. During template instantiation, or in a
+ /// non-dependent context, these are used to construct a SYCLKernelCallStmt.
+ ///
+ /// KernelInfoType is the type to be passed as the explicit template argument
+ /// to SYCL runtime libraries.
+ ///
+ /// KernelLaunchIdExpr is an UnresolvedLookupExpr or UnresolvedMemberExpr
+ /// for the SYCL kernel launch function, with KernelInfoType used as the
+ /// explicit template argument.
+ struct SYCLKernelCallStmtASTFragments {
+ QualType KernelInfoType;
+ Expr *KernelLaunchIdExpr;
+ };
+
+ /// BuildSYCLKernelCallStmtASTFragments builds portions of the AST needed
+ /// to construct SYCLKernelCallStmt or UnresolvedSYCLKernelCallStmt that
+ /// must be constructed early in the definition of a function before the
+ /// body of the function has been parsed.
+ ///
+ /// 'FD' must be a function declared with a valid sycl_kernel_entry_point
+ /// attribute and the current Sema context must match 'FD'.
+ ///
+ /// If construction of the AST fragments fails, diagnostics are issued and
+ /// a disengaged optional value is returned.
+ std::optional<SYCLKernelCallStmtASTFragments>
+ BuildSYCLKernelCallStmtASTFragments(FunctionDecl *FD);
+
+ /// BuildSYCLKernelCallStmt constructs a SYCLKernelCallStmt or an
+ /// UnresolvedSYCLKernelCallStmt depending on whether 'FD' is a templated
+ /// function.
+ ///
+ /// 'FD' must be a function declared with a valid sycl_kernel_entry_point
+ /// attribute, must not have had its body assigned yet, and the current
+ /// Sema context must match 'FD'.
+ ///
+ /// 'Body' is the parsed compound statement body of 'FD' to be wrapped by
+ /// the new SYCLKernelCallStmt or UnresolvedSYCLKernelCallStmt.
+ ///
+ /// 'ASTFragments' contains the AST fragments previously returned by a call
+ /// to BuildSYCLKernelCallStmtASTFragments() for 'FD'.
+ StmtResult
+ BuildSYCLKernelCallStmt(FunctionDecl *FD, CompoundStmt *Body,
+ const SYCLKernelCallStmtASTFragments &ASTFragments);
};
} // namespace clang
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index cb876cf8dd936..1d5ccf1c793eb 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -16513,27 +16513,23 @@ Decl *Sema::ActOnStartOfFunctionDef(Scope *FnBodyScope, Decl *D,
if (!FD->isInvalidDecl() && FD->hasAttr<SYCLKernelEntryPointAttr>() &&
FnBodyScope) {
- // An implicit call expression is synthesized for functions declared with
- // the sycl_kernel_entry_point attribute. The call may resolve to a
- // function template, a member function template, or a call operator
- // of a variable template depending on the results of unqualified lookup
- // for 'sycl_kernel_launch' from the beginning of the function body.
- // Performing that lookup requires the stack of parsing scopes active
- // when the definition is parsed and is thus done here; the result is
- // cached in FunctionScopeInfo and used to synthesize the (possibly
- // unresolved) call expression after the function body has been parsed.
+ // A function declared with the sycl_kernel_entry_point attribute has its
+ // body wrapped in a SYCLKernelCallStmt or UnresolvedSYCLKernelCallStmt.
+ // Those statements cannot be constructed until the function body has been
+ // parsed, but portions of the AST needed for construction have to be
+ // constructed early, before the function body has been parsed. Those are
+ // collected here.
const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>();
if (!SKEPAttr->isInvalidAttr()) {
- ExprResult LaunchIdExpr =
- SYCL().BuildSYCLKernelLaunchIdExpr(FD, SKEPAttr->getKernelName());
- // Do not mark 'FD' as invalid if construction of `LaunchIDExpr` produces
- // an invalid result. Name lookup failure for 'sycl_kernel_launch' is
- // treated as an error in the definition of 'FD'; treating it as an error
- // of the declaration would affect overload resolution which would
- // potentially result in additional errors. If construction of
- // 'LaunchIDExpr' failed, then 'SYCLKernelLaunchIdExpr' will be assigned
- // a null pointer value below; that is expected.
- getCurFunction()->SYCLKernelLaunchIdExpr = LaunchIdExpr.get();
+ // If construction of the AST fragments fails, SYCLKernelASTFragments
+ // will be assigned a disengaged value. This is expected and checked
+ // after the function body is parsed. Failure to construct the AST
+ // fragments is treated as an error in the definition of 'FD'; 'FD'
+ // is not marked as invalid since treating such failures as an error
+ // in the declaration would affect overload resolution which would
+ // potentially result in additional errors.
+ getCurFunction()->SYCLKernelASTFragments =
+ SYCL().BuildSYCLKernelCallStmtASTFragments(FD);
}
}
@@ -16738,30 +16734,26 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body, bool IsInstantiation,
SKEPAttr->setInvalidAttr();
}
- // Build an unresolved SYCL kernel call statement for a function template,
- // validate that a SYCL kernel call statement was instantiated for an
- // (implicit or explicit) instantiation of a function template, or otherwise
- // build a (resolved) SYCL kernel call statement for a non-templated
- // function or an explicit specialization.
+ // Build a SYCL kernel call statement to wrap the function body.
if (Body && !SKEPAttr->isInvalidAttr()) {
StmtResult SR;
if (FD->isTemplateInstantiation()) {
// The function body should already be a SYCLKernelCallStmt in this
// case, but might not be if there were previous errors.
SR = Body;
- } else if (!getCurFunction()->SYCLKernelLaunchIdExpr) {
- // If name lookup for a template named sycl_kernel_launch failed
- // earlier, don't try to build a SYCL kernel call statement as that
- // would cause additional errors to be issued; just proceed with the
- // original function body.
+ } else if (!getCurFunction()->SYCLKernelASTFragments) {
+ // Construction of the AST fragments failed earlier and diagnostics
+ // will have already been issued. Don't try to build a SYCL kernel
+ // call statement as that would cause additional errors to be issued;
+ // proceed with the original function body.
SR = Body;
- } else if (FD->isTemplated()) {
- SR = SYCL().BuildUnresolvedSYCLKernelCallStmt(
- cast<CompoundStmt>(Body), getCurFunction()->SYCLKernelLaunchIdExpr);
} else {
+ // Construct the SYCL kernel call statement. This will build an
+ // UnresolvedSYCLKernelCallStmt if FD is templated and a
+ // SYCLKernelCallStmt otherwise.
SR = SYCL().BuildSYCLKernelCallStmt(
FD, cast<CompoundStmt>(Body),
- getCurFunction()->SYCLKernelLaunchIdExpr);
+ *getCurFunction()->SYCLKernelASTFragments);
}
// If construction of the replacement body fails, just continue with the
// original function body. An early error return here is not valid; the
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index b942f19761f40..4de94621a5f19 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -425,8 +425,204 @@ void SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD) {
}
}
-ExprResult SemaSYCL::BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD,
- QualType KNT) {
+ClassTemplateDecl *SemaSYCL::GetSYCLKernelInfoClassTemplate() {
+ if (!SYCLKernelInfoClassTemplate) {
+ ASTContext &Ctx = getASTContext();
+ DeclContext *DC = Ctx.getTranslationUnitDecl();
+ SourceLocation Loc;
+
+ // Build a reserved identifier for the name of the class template. The
+ // name is formally unspecified.
+ IdentifierInfo *KernelInfoID =
+ &Ctx.Idents.get("__sycl_kernel_info", tok::TokenKind::identifier);
+
+ // Build a template parameter list consisting of a single type template
+ // parameter corresponding to a SYCL kernel name type. A reserved
+ // identifier is used for the formally unspecified name of the template
+ // parameter.
+ IdentifierInfo *KernelNameID =
+ &Ctx.Idents.get("__sycl_kernel_name", tok::TokenKind::identifier);
+ TemplateTypeParmDecl *KernelNameParam = TemplateTypeParmDecl::Create(
+ Ctx, DC, /*KeyLoc*/ Loc, /*NameLoc*/ Loc, 0, 0, KernelNameID,
+ /*Typename*/ true, /*ParameterPack*/ false);
+ TemplateParameterList *ParamList = TemplateParameterList::Create(
+ Ctx, /*TemplateLoc*/ Loc, /*LAngleLoc*/ Loc, KernelNameParam,
+ /*RAngleLoc*/ Loc, /*RequiresClause*/ nullptr);
+
+ // Build the class template and register it with its corresponding
+ // DeclContext.
+ CXXRecordDecl *KernelInfoDecl =
+ CXXRecordDecl::Create(Ctx, TagTypeKind::Class, DC, /*StartLoc*/ Loc,
+ /*IdLoc*/ Loc, KernelInfoID);
+ SYCLKernelInfoClassTemplate = ClassTemplateDecl::Create(
+ Ctx, DC, /*L*/ Loc, DeclarationName(KernelInfoID), ParamList,
+ KernelInfoDecl);
+ KernelInfoDecl->setDescribedClassTemplate(SYCLKernelInfoClassTemplate);
+ SYCLKernelInfoClassTemplate->setAccess(AS_public);
+ DC->addDecl(SYCLKernelInfoClassTemplate);
+ }
+ return SYCLKernelInfoClassTemplate;
+}
+
+QualType SemaSYCL::GetSYCLKernelInfoClassSpecializationType(QualType KNT) {
+ ASTContext &Ctx = getASTContext();
+
+ // Retrieve the class template for which a specialization is requested.
+ ClassTemplateDecl *KernelInfoTemplateDecl = GetSYCLKernelInfoClassTemplate();
+
+ // Build the template argument for the requested specialization.
+ TemplateArgument KNTA = TemplateArgument(KNT);
+
+ // If KNT is a dependent type, then a canonical type without an underlying
+ // type is returned. Otherwise, an explicit specialization declaration for
+ // a canonical underlying type is found or created and a canonical type
+ // returned for it.
+ QualType UnderlyingKernelInfoType;
+ if (!KNT->isDependentType()) {
+ // KNT is not a dependent type, check for a previously synthesized explicit
+ // specialization declaration.
+ CanQualType CanonicalKNT = Ctx.getCanonicalType(KNT);
+ auto IT = SYCLKernelInfoClassTemplateSpecializations.find(CanonicalKNT);
+ if (IT == SYCLKernelInfoClassTemplateSpecializations.end()) {
+ // No previous explicit specialization declaration found; synthesize one.
+ // KNT and KNTA are reassigned to avoid inadvertent use of a non-canonical
+ // type.
+ KNT = CanonicalKNT;
+ KNTA = TemplateArgument(KNT);
+
+ // Find the insertion position for the new specialization and ensure no
+ // such specialization already exists.
+ void *InsertPos;
+ ClassTemplateSpecializationDecl *KernelInfoDecl =
+ KernelInfoTemplateDecl->findSpecialization(KNTA, InsertPos);
+ assert(KernelInfoDecl == nullptr &&
+ "__sycl_kernel_info specialization unexpectedly already exists");
+
+ // Declare the explicit specialization, register it with the class
+ // template, and add it to the corresponding DeclContext.
+ DeclContext *DC = KernelInfoTemplateDecl->getDeclContext();
+ SourceLocation Loc;
+ KernelInfoDecl = ClassTemplateSpecializationDecl::Create(
+ Ctx, KernelInfoTemplateDecl->getTemplatedDecl()->getTagKind(), DC,
+ /*StartLoc*/ Loc, /*IdLoc*/ Loc, KernelInfoTemplateDecl, KNTA,
+ /*StrictPackMatch*/ false, /*PrevDecl*/ nullptr);
+ KernelInfoDecl->setSpecializationKind(TSK_ExplicitSpecialization);
+ KernelInfoTemplateDecl->AddSpecialization(KernelInfoDecl, InsertPos);
+ DC->addDecl(KernelInfoDecl);
+
+ // Cache the new declaration for subsequent calls.
+ SYCLKernelInfoClassTemplateSpecializations.insert(
+ std::make_pair(CanonicalKNT, KernelInfoDecl));
+
+ // Retrieve the canonical underlying type to use for the returned
+ // specialization type.
+ UnderlyingKernelInfoType = Ctx.getCanonicalTagType(KernelInfoDecl);
+ }
+ }
+
+ QualType KernelInfoType = Ctx.getTemplateSpecializationType(
+ ElaboratedTypeKeyword::None, TemplateName(KernelInfoTemplateDecl), KNTA,
+ /*CanonicalArgs*/ {}, UnderlyingKernelInfoType);
+
+ return KernelInfoType;
+}
+
+namespace {
+
+/// BuildSYCLKernelInfoClassSpecializationDefn builds an explicit
+/// specialization definition for the SYCL kernel info class template
+/// specialized for 'KNT'. 'KNT' cannot be a dependent type. The
+/// synthesized definition has the following structure where
+/// '<kernel-name>' represents the generated name of the offload
+/// kernel entry point function.
+///
+/// template<>
+/// struct sycl-kernel-info<KNT> {
+/// using kernel_name = KNT;
+/// static constexpr const char kernel_entry_point_name[] = "<kernel-name>";
+/// };
+ClassTemplateSpecializationDecl *
+BuildSYCLKernelInfoClassSpecializationDefn(Sema &SemaRef, QualType KNT) {
+ assert(!KNT->isDependentType());
+
+ ASTContext &Ctx = SemaRef.getASTContext();
+
+ // Retrieve the class template to be specialized.
+ ClassTemplateDecl *KernelInfoTemplateDecl =
+ SemaRef.SYCL().GetSYCLKernelInfoClassTemplate();
+
+ // Retrieve the specialization declaration that should have already been
+ // generated and validate that it corresponds to an explicit specialization
+ // as opposed to a somehow instantiated implicit specialization.
+ TemplateArgument KNTA = TemplateArgument(KNT);
+ void *InsertPos;
+ ClassTemplateSpecializationDecl *KernelInfoDecl =
+ KernelInfoTemplateDecl->findSpecialization(KNTA, InsertPos);
+ assert(KernelInfoDecl != nullptr &&
+ "__sycl_kernel_info specialization unexpectedly not found");
+ assert(KernelInfoDecl->getSpecializationKind() == TSK_ExplicitSpecialization);
+
+ // Start defining the specialization.
+ KernelInfoDecl->startDefinition();
+
+ // An injected class declaration would normally be declared for an explicit
+ // specialization, but isn't needed, so is skipped.
+
+ // Build the kernel_name type alias member.
+ SourceLocation Loc;
+ TypeSourceInfo *TSI = Ctx.getTrivialTypeSourceInfo(KNT, Loc);
+ IdentifierInfo &KernelNameTypeAliasID =
+ Ctx.Idents.get("kernel_name", tok::TokenKind::identifier);
+ TypeAliasDecl *KernelNameTypeAliasDecl = TypeAliasDecl::Create(
+ Ctx, KernelInfoDecl, Loc, Loc, &KernelNameTypeAliasID, TSI);
+ KernelNameTypeAliasDecl->setAccess(AS_public);
+ KernelInfoDecl->addDecl(KernelNameTypeAliasDecl);
+
+ // Build the kernel_entry_point_name static data member with its initializer.
+ const SYCLKernelInfo &SKI = Ctx.getSYCLKernelInfo(KNT);
+ const std::string &KernelName = SKI.GetKernelName();
+ QualType KernelNameCharTy = Ctx.CharTy;
+ QualType KernelNameArrayTy =
+ Ctx.getStringLiteralArrayType(KernelNameCharTy, KernelName.size());
+ IdentifierInfo &KernelEntryPointNameID =
+ Ctx.Idents.get("kernel_entry_point_name", tok::TokenKind::identifier);
+ VarDecl *KernelEntryPointNameDecl = VarDecl::Create(
+ Ctx, KernelInfoDecl, Loc, Loc, &KernelEntryPointNameID, KernelNameArrayTy,
+ Ctx.getTrivialTypeSourceInfo(KernelNameArrayTy, Loc), SC_Static);
+ Expr *KernelNameExpr =
+ StringLiteral::Create(Ctx, KernelName, StringLiteralKind::Ordinary,
+ /*Pascal*/ false, KernelNameArrayTy, Loc);
+ KernelEntryPointNameDecl->setConstexpr(true);
+ KernelEntryPointNameDecl->setImplicitlyInline();
+ KernelEntryPointNameDecl->setAccess(AS_public);
+ KernelEntryPointNameDecl->setInitStyle(VarDecl::CInit);
+ SemaRef.AddInitializerToDecl(KernelEntryPointNameDecl, KernelNameExpr,
+ /*DirectInit*/ false);
+ KernelInfoDecl->addDecl(KernelEntryPointNameDecl);
+
+ // Complete the definition of the specialization.
+ KernelInfoDecl->completeDefinition();
+
+ return KernelInfoDecl;
+}
+
+/// BuildSYCLKernelLaunchIdExpr performs lookup for a 'sycl_kernel_launch'
+/// template with 'KIT' used as an explicit type template argument. If
+/// lookup fails, a diagnostic is issued and an error expression returned.
+/// Otherwise, an UnresolvedLookupExpr, UnresolvedMemberExpr, DeclRefExpr,
+/// or MemberExpr will be returned. The resulting expression is intended
+/// to be passed as the 'IdExpr' argument to BuildSYCLKernelLaunchCallStmt()
+/// after the function body has been parsed.
+///
+/// 'FD' must be a function declared with a valid sycl_kernel_entry_point
+/// attribute and the current Sema context must match 'FD' so that scopes
+/// are properly in place such that lookup is performed as if from the first
+/// statement of the function body.
+///
+/// 'KIT' is the SYCL kernel info type to be used for lookup and it may
+/// be a dependent type.
+ExprResult BuildSYCLKernelLaunchIdExpr(Sema &SemaRef, FunctionDecl *FD,
+ QualType KIT) {
// The current context must be the function definition context to ensure
// that name lookup is performed within the correct scope.
assert(SemaRef.CurContext == FD && "The current declaration context does not "
@@ -466,9 +662,9 @@ ExprResult SemaSYCL::BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD,
return ExprError();
TemplateArgumentListInfo TALI{Loc, Loc};
- TemplateArgument KNTA = TemplateArgument(KNT);
+ TemplateArgument KITA = TemplateArgument(KIT);
TemplateArgumentLoc TAL =
- SemaRef.getTrivialTemplateArgumentLoc(KNTA, QualType(), Loc);
+ SemaRef.getTrivialTemplateArgumentLoc(KITA, QualType(), Loc);
TALI.addArgument(TAL);
ExprResult IdExpr;
@@ -493,14 +689,11 @@ ExprResult SemaSYCL::BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD,
return IdExpr;
}
-namespace {
-
// Constructs the arguments to be passed for the SYCL kernel launch call.
// The first argument is a string literal that contains the SYCL kernel
// name. The remaining arguments are the parameters of 'FD' passed as
// move-elligible xvalues. Returns true on error and false otherwise.
bool BuildSYCLKernelLaunchCallArgs(Sema &SemaRef, FunctionDecl *FD,
- const SYCLKernelInfo *SKI,
SmallVectorImpl<Expr *> &Args,
SourceLocation Loc) {
// The current context must be the function definition context to ensure
@@ -508,19 +701,6 @@ bool BuildSYCLKernelLaunchCallArgs(Sema &SemaRef, FunctionDecl *FD,
assert(SemaRef.CurContext == FD && "The current declaration context does not "
"match the requested function context");
- // Prepare a string literal that contains the kernel name.
- ASTContext &Ctx = SemaRef.getASTContext();
- 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);
-
// Forward all parameters of 'FD' to the SYCL kernel launch function as if
// by std::move().
for (ParmVarDecl *PVD : FD->parameters()) {
@@ -542,8 +722,8 @@ bool BuildSYCLKernelLaunchCallArgs(Sema &SemaRef, FunctionDecl *FD,
// Constructs the SYCL kernel launch call.
StmtResult BuildSYCLKernelLaunchCallStmt(Sema &SemaRef, FunctionDecl *FD,
- const SYCLKernelInfo *SKI,
- Expr *IdExpr, SourceLocation Loc) {
+ QualType InfoDeclType, Expr *IdExpr,
+ SourceLocation Loc) {
SmallVector<Stmt *> Stmts;
// IdExpr may be null if name lookup failed.
if (IdExpr) {
@@ -557,7 +737,7 @@ StmtResult BuildSYCLKernelLaunchCallStmt(Sema &SemaRef, FunctionDecl *FD,
CSC.Entity = FD;
Sema::ScopedCodeSynthesisContext ScopedCSC(SemaRef, CSC);
- if (BuildSYCLKernelLaunchCallArgs(SemaRef, FD, SKI, Args, Loc))
+ if (BuildSYCLKernelLaunchCallArgs(SemaRef, FD, Args, Loc))
return StmtError();
}
@@ -769,19 +949,26 @@ bool verifyKernelParams(FunctionDecl *FD, SemaSYCL &SemaSYCLRef) {
return KAC.isInvalid();
}
-} // unnamed namespace
-
-StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl *FD,
- CompoundStmt *Body,
- Expr *LaunchIdExpr) {
+/// BuildResolvedSYCLKernelCallStmt builds a SYCLKernelCallStmt to serve as
+/// the body of 'FD'. If construction fails, a diagnostic is issued and an
+/// error statement is returned.
+///
+/// 'FD' must be a function declared with a valid sycl_kernel_entry_point
+/// attribute and must not have had its body assigned yet.
+///
+/// 'Body' is the parsed compound statement body of 'FD' to be wrapped by
+/// the new SYCLKernelCallStmt.
+///
+/// 'ASTFragments' contains the AST fragments previously returned by a call
+/// to BuildSYCLKernelCallStmtASTFragments() for 'FD'.
+StmtResult BuildResolvedSYCLKernelCallStmt(
+ Sema &SemaRef, FunctionDecl *FD, CompoundStmt *Body,
+ const SemaSYCL::SYCLKernelCallStmtASTFragments &ASTFragments) {
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.
- assert(SemaRef.CurContext == FD && "The current declaration context does not "
- "match the requested function context");
+
+ ASTContext &Ctx = SemaRef.getASTContext();
const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>();
assert(SKEPAttr && "Missing sycl_kernel_entry_point attribute");
@@ -790,13 +977,17 @@ StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl *FD,
// Ensure that the kernel name was previously registered and that the
// stored declaration matches.
- const SYCLKernelInfo &SKI =
- getASTContext().getSYCLKernelInfo(SKEPAttr->getKernelName());
+ const SYCLKernelInfo &SKI = Ctx.getSYCLKernelInfo(SKEPAttr->getKernelName());
assert(declaresSameEntity(SKI.getKernelEntryPointDecl(), FD) &&
"SYCL kernel name conflict");
- if (verifyKernelParams(FD, *this))
+
+ if (verifyKernelParams(FD, SemaRef.SYCL()))
return StmtError();
+ // Inject an explicit specialization declaration for a specialization
+ // of __sycl_kernel_info parameterized by the kernel name type.
+ BuildSYCLKernelInfoClassSpecializationDefn(SemaRef, SKI.getKernelNameType());
+
// Build the outline of the synthesized device entry point function.
OutlinedFunctionDecl *OFD =
BuildSYCLKernelEntryPointOutline(SemaRef, FD, Body);
@@ -806,18 +997,70 @@ StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl *FD,
// is required to emit diagnostics.
SourceLocation Loc = Body->getLBracLoc();
StmtResult LaunchResult =
- BuildSYCLKernelLaunchCallStmt(SemaRef, FD, &SKI, LaunchIdExpr, Loc);
+ BuildSYCLKernelLaunchCallStmt(SemaRef, FD, ASTFragments.KernelInfoType,
+ ASTFragments.KernelLaunchIdExpr, Loc);
if (LaunchResult.isInvalid())
return StmtError();
- Stmt *NewBody =
- new (getASTContext()) SYCLKernelCallStmt(Body, LaunchResult.get(), OFD);
+ Stmt *NewBody = new (Ctx) SYCLKernelCallStmt(Body, LaunchResult.get(), OFD);
return NewBody;
}
-StmtResult SemaSYCL::BuildUnresolvedSYCLKernelCallStmt(CompoundStmt *Body,
- Expr *LaunchIdExpr) {
+/// BuildUnresolvedSYCLKernelCallStmt constructs a UnresolvedSYCLKernelCallStmt
+/// to serve as the body of a function in a dependent context.
+///
+/// 'Body' is the parsed compound statement body to be wrapped by the new
+/// UnresolvedSYCLKernelCallStmt.
+///
+/// 'ASTFragments' contains the AST fragments previously returned by a call
+/// to BuildSYCLKernelCallStmtASTFragments() for the function.
+StmtResult BuildUnresolvedSYCLKernelCallStmt(
+ Sema &SemaRef, CompoundStmt *Body,
+ const SemaSYCL::SYCLKernelCallStmtASTFragments &ASTFragments) {
return UnresolvedSYCLKernelCallStmt::Create(SemaRef.getASTContext(), Body,
- LaunchIdExpr);
+ ASTFragments.KernelInfoType,
+ ASTFragments.KernelLaunchIdExpr);
+}
+
+} // unnamed namespace
+
+std::optional<SemaSYCL::SYCLKernelCallStmtASTFragments>
+SemaSYCL::BuildSYCLKernelCallStmtASTFragments(FunctionDecl *FD) {
+ const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>();
+ assert(SKEPAttr && "FD doesn't have a sycl_kernel_entry_point attribute");
+ assert(!SKEPAttr->isInvalidAttr() &&
+ "FD has an invalid sycl_kernel_entry_point attribute");
+ QualType KernelName = SKEPAttr->getKernelName();
+
+ // Retrieve or build the kernel info class template specialization type
+ // needed for lookup of SYCL runtime declared templates.
+ QualType KernelInfoType =
+ GetSYCLKernelInfoClassSpecializationType(KernelName);
+ assert(!KernelInfoType.isNull());
+
+ // Perform lookup for the sycl_kernel_launch template.
+ ExprResult KernelLaunchIdER =
+ BuildSYCLKernelLaunchIdExpr(SemaRef, FD, KernelInfoType);
+ if (KernelLaunchIdER.isInvalid())
+ return {};
+
+ return SemaSYCL::SYCLKernelCallStmtASTFragments{KernelInfoType,
+ KernelLaunchIdER.get()};
+}
+
+StmtResult SemaSYCL::BuildSYCLKernelCallStmt(
+ FunctionDecl *FD, CompoundStmt *Body,
+ const SemaSYCL::SYCLKernelCallStmtASTFragments &ASTFragments) {
+ // 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 && "The current declaration context does not "
+ "match the requested function context");
+
+ if (FD->isTemplated()) {
+ return BuildUnresolvedSYCLKernelCallStmt(SemaRef, Body, ASTFragments);
+ } else {
+ return BuildResolvedSYCLKernelCallStmt(SemaRef, FD, Body, ASTFragments);
+ }
}
diff --git a/clang/lib/Sema/SemaTemplateInstantiate.cpp b/clang/lib/Sema/SemaTemplateInstantiate.cpp
index 990e52eb27d39..9e2f9eee2ad3c 100644
--- a/clang/lib/Sema/SemaTemplateInstantiate.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp
@@ -1275,7 +1275,8 @@ void Sema::PrintInstantiationStack(InstantiationContextDiagFuncRef DiagFunc) {
DiagFunc(SKEPAttr->getLocation(), PDiag(diag::note_sycl_runtime_defect));
DiagFunc(SKEPAttr->getLocation(),
PDiag(diag::note_sycl_kernel_launch_lookup_here)
- << SKEPAttr->getKernelName());
+ << SYCL().GetSYCLKernelInfoClassSpecializationType(
+ SKEPAttr->getKernelName()));
break;
}
case CodeSynthesisContext::SYCLKernelLaunchOverloadResolution: {
@@ -1287,7 +1288,8 @@ void Sema::PrintInstantiationStack(InstantiationContextDiagFuncRef DiagFunc) {
DiagFunc(SKEPAttr->getLocation(), PDiag(diag::note_sycl_runtime_defect));
DiagFunc(SKEPAttr->getLocation(),
PDiag(diag::note_sycl_kernel_launch_overload_resolution_here)
- << SKEPAttr->getKernelName()
+ << SYCL().GetSYCLKernelInfoClassSpecializationType(
+ SKEPAttr->getKernelName())
<< convertCallArgsValueCategoryAndTypeToString(
*this, llvm::ArrayRef(Active->CallArgs,
Active->NumCallArgs)));
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 3c8fcbe582b43..1a301d82603c0 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -13173,6 +13173,15 @@ StmtResult TreeTransform<Derived>::TransformUnresolvedSYCLKernelCallStmt(
if (!SKEPAttr || SKEPAttr->isInvalidAttr())
return StmtError();
+ // The kernel name type held by the transformed SKEPAttr should no longer be
+ // dependent. Ensure that is so and retrieve a non-dependent kernel info
+ // specialization type.
+ assert(!SKEPAttr->getKernelName()->isDependentType() &&
+ "SYCL kernel name type should not be dependent here");
+ QualType InfoType = SemaRef.SYCL().GetSYCLKernelInfoClassSpecializationType(
+ SKEPAttr->getKernelName());
+ assert(!InfoType.isNull());
+
ExprResult IdExpr = getDerived().TransformExpr(S->getKernelLaunchIdExpr());
if (IdExpr.isInvalid())
return StmtError();
@@ -13181,9 +13190,9 @@ StmtResult TreeTransform<Derived>::TransformUnresolvedSYCLKernelCallStmt(
if (Body.isInvalid())
return StmtError();
+ SemaSYCL::SYCLKernelCallStmtASTFragments ASTFragments{InfoType, IdExpr.get()};
StmtResult SR = SemaRef.SYCL().BuildSYCLKernelCallStmt(
- cast<FunctionDecl>(SemaRef.CurContext), cast<CompoundStmt>(Body.get()),
- IdExpr.get());
+ FD, cast<CompoundStmt>(Body.get()), ASTFragments);
if (SR.isInvalid())
return StmtError();
diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp
index 09e31a9732cdf..27b7c29ce16a2 100644
--- a/clang/lib/Serialization/ASTReaderStmt.cpp
+++ b/clang/lib/Serialization/ASTReaderStmt.cpp
@@ -618,6 +618,7 @@ void ASTStmtReader::VisitUnresolvedSYCLKernelCallStmt(
VisitStmt(S);
S->setOriginalStmt(cast<CompoundStmt>(Record.readSubStmt()));
+ S->setKernelInfoType(Record.readType());
S->setKernelLaunchIdExpr(Record.readExpr());
}
diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp
index c193e9d0ea1fe..49f0ac57bab8e 100644
--- a/clang/lib/Serialization/ASTWriterStmt.cpp
+++ b/clang/lib/Serialization/ASTWriterStmt.cpp
@@ -702,6 +702,7 @@ void ASTStmtWriter::VisitUnresolvedSYCLKernelCallStmt(
VisitStmt(S);
Record.AddStmt(S->getOriginalStmt());
+ Record.AddTypeRef(S->getKernelInfoType());
Record.AddStmt(S->getKernelLaunchIdExpr());
Code = serialization::STMT_UNRESOLVED_SYCL_KERNEL_CALL;
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 d66d4fdcc9483..aaccb320a9cf2 100644
--- a/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp
+++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp
@@ -34,8 +34,8 @@ template<int> struct K {
void operator()(Ts...) const {}
};
-template<typename KernelName, typename... Ts>
-void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KernelInfo, typename... Ts>
+void sycl_kernel_launch(Ts...) {}
[[clang::sycl_kernel_entry_point(KN<1>)]]
void skep1() {
@@ -45,10 +45,8 @@ void skep1() {
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
// 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: | | | `-ImplicitCastExpr {{.*}} 'void (*)()' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'void ()' lvalue Function {{.*}} 'sycl_kernel_launch' {{.*}}
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
// CHECK-NEXT: | | `-CompoundStmt {{.*}}
// CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} KN<1>
@@ -69,10 +67,20 @@ void skep2<KN<2>>(K<2>);
// CHECK-NEXT: | | | |-CompoundStmt {{.*}}
// CHECK-NEXT: | | | | `-CallExpr {{.*}} '<dependent type>'
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT'
+// CHECK-NEXT: | | | |-TemplateSpecializationType {{.*}} '__sycl_kernel_info<KNT>' dependent
+// CHECK-NEXT: | | | | |-name: '__sycl_kernel_info'
+// CHECK-NEXT: | | | | | `-ClassTemplateDecl {{.*}} referenced __sycl_kernel_info external-linkage
+// CHECK-NEXT: | | | | `-TemplateArgument type 'KNT':'type-parameter-0-0'
+// CHECK-NEXT: | | | | `-TemplateTypeParmType {{.*}} 'KNT' dependent depth 0 index 0
+// CHECK-NEXT: | | | | `-TemplateTypeParm {{.*}} 'KNT'
// CHECK-NEXT: | | | `-UnresolvedLookupExpr {{.*}} '<dependent type>' lvalue (ADL) = 'sycl_kernel_launch' {{.*}}
-// CHECK-NEXT: | | | `-TemplateArgument type 'KNT':'type-parameter-0-0'
-// CHECK-NEXT: | | | `-TemplateTypeParmType {{.*}} 'KNT' dependent depth 0 index 0
-// CHECK-NEXT: | | | `-TemplateTypeParm {{.*}} 'KNT'
+// CHECK-NEXT: | | | `-TemplateArgument type '__sycl_kernel_info<KNT>':'__sycl_kernel_info<type-parameter-0-0>'
+// CHECK-NEXT: | | | `-TemplateSpecializationType {{.*}} '__sycl_kernel_info<KNT>' dependent
+// CHECK-NEXT: | | | |-name: '__sycl_kernel_info'
+// CHECK-NEXT: | | | | `-ClassTemplateDecl {{.*}} referenced __sycl_kernel_info external-linkage
+// CHECK-NEXT: | | | `-TemplateArgument type 'KNT':'type-parameter-0-0'
+// CHECK-NEXT: | | | `-TemplateTypeParmType {{.*}} 'KNT' dependent depth 0 index 0
+// CHECK-NEXT: | | | `-TemplateTypeParm {{.*}} 'KNT'
// CHECK-NEXT: | | `-SYCLKernelEntryPointAttr {{.*}} KNT
// CHECK-NEXT: | `-FunctionDecl {{.*}} skep2 'void (K<2>)' explicit_instantiation_definition instantiated_from 0x{{.+}} external-linkage
@@ -92,10 +100,8 @@ void skep2<KN<2>>(K<2>);
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'K<2>' lvalue ParmVar {{.*}} 'k' 'K<2>'
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
// 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: | | | |-ImplicitCastExpr {{.*}} 'void (*)(K<2>)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (K<2>)' lvalue Function {{.*}} 'sycl_kernel_launch' {{.*}}
// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'K<2>' 'void (K<2> &&) noexcept'
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'K<2>' xvalue <NoOp>
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'K<2>' lvalue ParmVar {{.*}} 'k' 'K<2>'
@@ -128,10 +134,20 @@ void skep3<KN<3>>(K<3> k) {
// CHECK-NEXT: | | | |-CompoundStmt {{.*}}
// CHECK-NEXT: | | | | `-CallExpr {{.*}} '<dependent type>'
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT'
+// CHECK-NEXT: | | | |-TemplateSpecializationType {{.*}} '__sycl_kernel_info<KNT>' dependent
+// CHECK-NEXT: | | | | |-name: '__sycl_kernel_info'
+// CHECK-NEXT: | | | | | `-ClassTemplateDecl {{.*}} referenced __sycl_kernel_info external-linkage
+// CHECK-NEXT: | | | | `-TemplateArgument type 'KNT':'type-parameter-0-0'
+// CHECK-NEXT: | | | | `-TemplateTypeParmType {{.*}} 'KNT' dependent depth 0 index 0
+// CHECK-NEXT: | | | | `-TemplateTypeParm {{.*}} 'KNT'
// CHECK-NEXT: | | | `-UnresolvedLookupExpr {{.*}} '<dependent type>' lvalue (ADL) = 'sycl_kernel_launch' {{.*}}
-// CHECK-NEXT: | | | `-TemplateArgument type 'KNT':'type-parameter-0-0'
-// CHECK-NEXT: | | | `-TemplateTypeParmType {{.*}} 'KNT' dependent depth 0 index 0
-// CHECK-NEXT: | | | `-TemplateTypeParm {{.*}} 'KNT'
+// CHECK-NEXT: | | | `-TemplateArgument type '__sycl_kernel_info<KNT>':'__sycl_kernel_info<type-parameter-0-0>'
+// CHECK-NEXT: | | | `-TemplateSpecializationType {{.*}} '__sycl_kernel_info<KNT>' dependent
+// CHECK-NEXT: | | | |-name: '__sycl_kernel_info'
+// CHECK-NEXT: | | | | `-ClassTemplateDecl {{.*}} referenced __sycl_kernel_info external-linkage
+// CHECK-NEXT: | | | `-TemplateArgument type 'KNT':'type-parameter-0-0'
+// CHECK-NEXT: | | | `-TemplateTypeParmType {{.*}} 'KNT' dependent depth 0 index 0
+// CHECK-NEXT: | | | `-TemplateTypeParm {{.*}} 'KNT'
// CHECK-NEXT: | | `-SYCLKernelEntryPointAttr {{.*}} KNT
// CHECK-NEXT: | `-Function {{.*}} 'skep3' 'void (K<3>)'
@@ -151,11 +167,9 @@ 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: | | | `-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: | | | `-CallExpr {{.*}} 'void'
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(K<3>)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (K<3>)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (K<3>)' {{.*}}
// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'K<3>' 'void (K<3> &&) noexcept'
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'K<3>' xvalue <NoOp>
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'K<3>' lvalue ParmVar {{.*}} 'k' 'K<3>'
@@ -190,10 +204,8 @@ void skep4(K<4> k, int p1, int p2) {
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'p2' 'int'
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
// 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: | | | |-ImplicitCastExpr {{.*}} 'void (*)(K<4>, int, int)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (K<4>, int, int)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (K<4>, int, int)' {{.*}}
// CHECK-NEXT: | | | |-CXXConstructExpr {{.*}} 'K<4>' 'void (K<4> &&) noexcept'
// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'K<4>' xvalue <NoOp>
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'K<4>' lvalue ParmVar {{.*}} 'k' 'K<4>'
@@ -235,10 +247,8 @@ void skep5(int unused1, K<5> k, int unused2, int p, int unused3) {
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
// CHECK: | | |-CompoundStmt {{.*}}
// 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 {{.*}} 'void (*)(int, K<5>, int, int, int)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (int, K<5>, int, int, int)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (int, K<5>, int, int, int)' {{.*}}
// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'int' xvalue <NoOp>
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} 'unused1' 'int'
@@ -301,10 +311,8 @@ void skep6(const S6 &k) {
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'const S6' lvalue ParmVar {{.*}} 'k' 'const S6 &'
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
// 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: | | | |-ImplicitCastExpr {{.*}} 'void (*)(S6)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (S6)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (S6)' {{.*}}
// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'S6' 'void (const S6 &) noexcept'
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'const S6' lvalue ParmVar {{.*}} 'k' 'const S6 &'
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
@@ -342,10 +350,8 @@ void skep7(S7 k) {
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'S7' lvalue ParmVar {{.*}} 'k' 'S7'
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
// 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: | | | |-ImplicitCastExpr {{.*}} 'void (*)(S7)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (S7)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (S7)' {{.*}}
// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'S7' 'void (S7 &&) noexcept'
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'S7' xvalue <NoOp>
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'S7' lvalue ParmVar {{.*}} 'k' 'S7'
@@ -376,10 +382,8 @@ void skep8(S8 k) {
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
// CHECK: | | |-CompoundStmt {{.*}}
// 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: | | | |-ImplicitCastExpr {{.*}} 'void (*)(S8)' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (S8)' lvalue Function {{.*}} 'sycl_kernel_launch' 'void (S8)' {{.*}}
// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'S8' 'void (S8 &&) noexcept'
// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'S8' xvalue <NoOp>
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'S8' lvalue ParmVar {{.*}} 'k' 'S8'
@@ -387,8 +391,8 @@ void skep8(S8 k) {
// CHECK: | `-SYCLKernelEntryPointAttr {{.*}}
class Handler {
- template <typename KNT, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...) {}
+ template <typename KIT, typename... Ts>
+ void sycl_kernel_launch(Ts...) {}
public:
template<typename KNT, typename KT>
[[clang::sycl_kernel_entry_point(KNT)]]
@@ -414,6 +418,12 @@ void foo() {
// 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: | | | | |-TemplateSpecializationType {{.*}} '__sycl_kernel_info<KNT>' dependent
+// CHECK-NEXT: | | | | | |-name: '__sycl_kernel_info'
+// CHECK-NEXT: | | | | | | `-ClassTemplateDecl {{.*}} <invalid sloc> referenced __sycl_kernel_info external-linkage
+// CHECK-NEXT: | | | | | `-TemplateArgument type 'KNT':'type-parameter-0-0'
+// CHECK-NEXT: | | | | | `-TemplateTypeParmType {{.*}} 'KNT' dependent depth 0 index 0
+// CHECK-NEXT: | | | | | `-TemplateTypeParm {{.*}} 'KNT'
// CHECK-NEXT: | | | | `-UnresolvedMemberExpr {{.*}} '<bound member function type>' lvalue
// CHECK-NEXT: | | | `-SYCLKernelEntryPointAttr {{.*}} KNT
// CHECK-NEXT: | | `-CXXMethodDecl {{.*}} used skep9 {{.*}} implicit_instantiation implicit-inline instantiated_from 0x{{.*}}
@@ -441,8 +451,6 @@ void foo() {
// 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 {{.*}} xvalue <NoOp>
// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} lvalue ParmVar {{.*}} 'k' {{.*}}
@@ -470,4 +478,4 @@ void foo() {
void the_end() {}
-// CHECK: `-FunctionDecl {{.*}} the_end 'void ()' external-linkage
+// CHECK: |-FunctionDecl {{.*}} the_end 'void ()' external-linkage
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 cca10365e7686..d27bba4fa6955 100644
--- a/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
+++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
@@ -28,8 +28,8 @@
// A unique kernel name type is required for each declared kernel entry point.
template<int, int=0> struct KN;
-template<typename KernelName, typename... Ts>
-void sycl_kernel_launch(const char *, Ts... Args) {}
+template<typename KernelInfo, typename... Ts>
+void sycl_kernel_launch(Ts... Args) {}
[[clang::sycl_kernel_entry_point(KN<1>)]]
void skep1() {
@@ -157,4 +157,4 @@ void skep7() {
// CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} KN<7>
void the_end() {}
-// CHECK: `-FunctionDecl {{.*}} the_end 'void ()' external-linkage
+// CHECK: |-FunctionDecl {{.*}} the_end 'void ()' external-linkage
diff --git a/clang/test/ASTSYCL/ast-print-sycl-kernel-call.cpp b/clang/test/ASTSYCL/ast-print-sycl-kernel-call.cpp
index 5adaa367ed9c1..5f0695a24a632 100644
--- a/clang/test/ASTSYCL/ast-print-sycl-kernel-call.cpp
+++ b/clang/test/ASTSYCL/ast-print-sycl-kernel-call.cpp
@@ -2,8 +2,8 @@
// RUN: %clang_cc1 -fsycl-is-device -ast-print %s -o - | FileCheck %s
struct sycl_kernel_launcher {
- template<typename KernelName, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...) {}
+ template<typename KernelInfo, typename... Ts>
+ void sycl_kernel_launch(Ts...) {}
template<typename KernelName, typename KernelType>
[[clang::sycl_kernel_entry_point(KernelName)]]
diff --git a/clang/test/CodeGenSYCL/function-attrs.cpp b/clang/test/CodeGenSYCL/function-attrs.cpp
index 60d3cf10055ec..da9c2efd56b36 100644
--- a/clang/test/CodeGenSYCL/function-attrs.cpp
+++ b/clang/test/CodeGenSYCL/function-attrs.cpp
@@ -26,7 +26,7 @@ int foo() {
return 1;
}
-template <typename Name, typename... Ts>
+template <typename KernelInfo, typename... Ts>
void sycl_kernel_launch(Ts...) {}
template <typename Name, typename Func>
diff --git a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
index 56558cec82c56..af1dec78a60f2 100644
--- a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
+++ b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
@@ -30,8 +30,8 @@
// these functions are used to generate offload kernel entry points (SYCL kernel
// caller functions).
-template <typename KernelName, typename... Ts>
-void sycl_kernel_launch(const char *, Ts...) {}
+template <typename KernelInfo, typename... Ts>
+void sycl_kernel_launch(Ts...) {}
struct single_purpose_kernel_name;
struct single_purpose_kernel {
@@ -53,8 +53,8 @@ void kernel_single_task(KernelType kernelFunc) {
struct \u03b4\u03c4\u03c7; // Delta Tau Chi (δτχ)
class handler {
- template <typename KernelName, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...) {}
+ template <typename KernelInfo, typename... Ts>
+ void sycl_kernel_launch(Ts...) {}
public:
template <typename KernelName, typename KernelType>
[[clang::sycl_kernel_entry_point(KernelName)]]
@@ -140,28 +140,11 @@ 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: @.str.3 = private unnamed_addr constant [15 x i8] c"_ZTSZ4mainE2KN\00", align 1
-// CHECK-HOST-LINUX: @.str.4 = private unnamed_addr constant [26 x i8] c"_ZTS19ref_arg_kernel_name\00", align 1
-// CHECK-HOST-LINUX @.str.5 = private unnamed_addr constant [30 x i8] c"_ZTS23fwd_ref_arg_kernel_name\00", align 1
-// CHECK-HOST-LINUX: @.str.6 = private unnamed_addr constant [35 x i8] c"_ZTS28fwd_ref_arg_kernel_name_move\00", align 1
-// CHECK-HOST-LINUX: @.str.7 = private unnamed_addr constant [33 x i8] c"_ZTS26rvalue_ref_arg_kernel_name\00", align 1
-// CHECK-HOST-WINDOWS: @"??_C at _0CB@KFIJOMLB at _ZTS26single_purpose_kernel_name@" = linkonce_odr dso_local unnamed_addr constant [33 x i8] c"_ZTS26single_purpose_kernel_name\00", comdat
-// CHECK-HOST-WINDOWS: @"??_C at _0BC@NHCDOLAA at _ZTSZ4mainEUlT_E_?$AA@" = linkonce_odr dso_local unnamed_addr constant [18 x i8] c"_ZTSZ4mainEUlT_E_\00", comdat
-// CHECK-HOST-WINDOWS: @"??_C at _0M@BCGAEMBE at _ZTS6?N?$LE?O?$IE?O?$IH?$AA@" = linkonce_odr dso_local unnamed_addr constant [12 x i8] c"_ZTS6\CE\B4\CF\84\CF\87\00", comdat
-// CHECK-HOST-WINDOWS: @"??_C at _0P@DLGHPODL at _ZTSZ4mainE2KN?$AA@" = linkonce_odr dso_local unnamed_addr constant [15 x i8] c"_ZTSZ4mainE2KN\00", comdat
-// CHECK-HOST-WINDOWS: @"??_C at _0BK@PPDJPOBM at _ZTS19ref_arg_kernel_name?$AA@" = linkonce_odr dso_local unnamed_addr constant [26 x i8] c"_ZTS19ref_arg_kernel_name\00", comdat
-// CHECK-HOST-WINDOWS: @"??_C at _0BO@KEIBIHKH at _ZTS23fwd_ref_arg_kernel_name?$AA@" = linkonce_odr dso_local unnamed_addr constant [30 x i8] c"_ZTS23fwd_ref_arg_kernel_name\00", comdat
-// CHECK-HOST-WINDOWS: @"??_C at _0CD@FDALJLMM at _ZTS28fwd_ref_arg_kernel_name_mo@" = linkonce_odr dso_local unnamed_addr constant [35 x i8] c"_ZTS28fwd_ref_arg_kernel_name_move\00", comdat
-// CHECK-HOST-WINDOWS: @"??_C at _0CB@HCPMABHM at _ZTS26rvalue_ref_arg_kernel_name@" = linkonce_odr dso_local unnamed_addr constant [33 x i8] c"_ZTS26rvalue_ref_arg_kernel_name\00", comdat
-//
// 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: %agg.tmp = alloca %struct.single_purpose_kernel, align 1
-// CHECK-HOST-LINUX-NEXT: call void @_Z18sycl_kernel_launchI26single_purpose_kernel_nameJ21single_purpose_kernelEEvPKcDpT0_(ptr noundef @.str)
+// CHECK-HOST-LINUX-NEXT: call void @_Z18sycl_kernel_launchI18__sycl_kernel_infoI26single_purpose_kernel_nameEJ21single_purpose_kernelEEvDpT0_()
// CHECK-HOST-LINUX-NEXT: ret void
// CHECK-HOST-LINUX-NEXT: }
//
@@ -174,7 +157,7 @@ int main() {
// 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_JS1_EEvPKcDpT0_(ptr noundef @.str.1, i32 %0)
+// CHECK-HOST-LINUX-NEXT: call void @_Z18sycl_kernel_launchI18__sycl_kernel_infoIZ4mainEUlT_E_EJS2_EEvDpT0_(i32 %0)
// CHECK-HOST-LINUX-NEXT: ret void
// CHECK-HOST-LINUX-NEXT: }
//
@@ -182,7 +165,7 @@ int main() {
// CHECK-HOST-LINUX-NEXT: entry:
// CHECK-HOST-LINUX-NEXT: %kernelFunc = alloca %class.anon.0, align 1
// 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\87JZ4mainEUliE_EEvPKcDpT0_"(ptr noundef @.str.2)
+// CHECK-HOST-LINUX-NEXT: call void @"_Z18sycl_kernel_launchI18__sycl_kernel_infoI6\CE\B4\CF\84\CF\87EJZ4mainEUliE_EEvDpT0_"()
// CHECK-HOST-LINUX-NEXT: ret void
// CHECK-HOST-LINUX-NEXT: }
@@ -202,7 +185,7 @@ int main() {
// CHECK-HOST-LINUX-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 %agg.tmp, ptr align 4 %k, i64 4, i1 false)
// 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_launchIZ4mainE2KNJZ4mainEUliiE_iiEEEvPKcDpT0_(ptr noundef nonnull align 1 dereferenceable(1) %this1, ptr noundef @.str.3, ptr noundef align 4 %agg.tmp, i32 noundef %0, i32 noundef %1)
+// CHECK-HOST-LINUX-NEXT: call void @_ZN7handler18sycl_kernel_launchI18__sycl_kernel_infoIZ4mainE2KNEJZ4mainEUliiE_iiEEEvDpT0_(ptr noundef nonnull align 1 dereferenceable(1) %this1, ptr noundef align 4 %agg.tmp, i32 noundef %0, i32 noundef %1)
// CHECK-HOST-LINUX-NEXT: call void @_ZZ4mainENUliiE_D1Ev(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) %agg.tmp) #{{[0-9]+}}
// CHECK-HOST-LINUX-NEXT: ret void
// CHECK-HOST-LINUX-NEXT: }
@@ -216,7 +199,7 @@ int main() {
// CHECK-HOST-LINUX-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 %agg.tmp, ptr align 4 %0, i64 4, i1 false)
// CHECK-HOST-LINUX-NEXT: %coerce.dive = getelementptr inbounds nuw %class.anon, ptr %agg.tmp, i32 0, i32 0
// CHECK-HOST-LINUX-NEXT: %1 = load i32, ptr %coerce.dive, align 4
-// CHECK-HOST-LINUX-NEXT: call void @_Z18sycl_kernel_launchI19ref_arg_kernel_nameJZ4mainEUlT_E_EEvPKcDpT0_(ptr noundef @.str.4, i32 %1)
+// CHECK-HOST-LINUX-NEXT: call void @_Z18sycl_kernel_launchI18__sycl_kernel_infoI19ref_arg_kernel_nameEJZ4mainEUlT_E_EEvDpT0_(i32 %1)
// CHECK-HOST-LINUX-NEXT: ret void
// CHECK-HOST-LINUX-NEXT: }
//
@@ -229,7 +212,7 @@ int main() {
// CHECK-HOST-LINUX-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 %agg.tmp, ptr align 4 %0, i64 4, i1 false)
// CHECK-HOST-LINUX-NEXT: %coerce.dive = getelementptr inbounds nuw %class.anon, ptr %agg.tmp, i32 0, i32 0
// CHECK-HOST-LINUX-NEXT: %1 = load i32, ptr %coerce.dive, align 4
-// CHECK-HOST-LINUX-NEXT: call void @_Z18sycl_kernel_launchI23fwd_ref_arg_kernel_nameJZ4mainEUlT_E_EEvPKcDpT0_(ptr noundef @.str.5, i32 %1)
+// CHECK-HOST-LINUX-NEXT: call void @_Z18sycl_kernel_launchI18__sycl_kernel_infoI23fwd_ref_arg_kernel_nameEJZ4mainEUlT_E_EEvDpT0_(i32 %1)
// CHECK-HOST-LINUX-NEXT: ret void
// CHECK-HOST-LINUX-NEXT: }
//
@@ -242,7 +225,7 @@ int main() {
// CHECK-HOST-LINUX-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 %agg.tmp, ptr align 4 %0, i64 4, i1 false)
// CHECK-HOST-LINUX-NEXT: %coerce.dive = getelementptr inbounds nuw %class.anon, ptr %agg.tmp, i32 0, i32 0
// CHECK-HOST-LINUX-NEXT: %1 = load i32, ptr %coerce.dive, align 4
-// CHECK-HOST-LINUX-NEXT: call void @_Z18sycl_kernel_launchI28fwd_ref_arg_kernel_name_moveJZ4mainEUlT_E_EEvPKcDpT0_(ptr noundef @.str.6, i32 %1)
+// CHECK-HOST-LINUX-NEXT: call void @_Z18sycl_kernel_launchI18__sycl_kernel_infoI28fwd_ref_arg_kernel_name_moveEJZ4mainEUlT_E_EEvDpT0_(i32 %1)
// CHECK-HOST-LINUX-NEXT: ret void
// CHECK-HOST-LINUX-NEXT: }
//
@@ -255,7 +238,7 @@ int main() {
// CHECK-HOST-LINUX-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 %agg.tmp, ptr align 4 %0, i64 4, i1 false)
// CHECK-HOST-LINUX-NEXT: %coerce.dive = getelementptr inbounds nuw %class.anon.2, ptr %agg.tmp, i32 0, i32 0
// CHECK-HOST-LINUX-NEXT: %1 = load i32, ptr %coerce.dive, align 4
-// CHECK-HOST-LINUX-NEXT: call void @_Z18sycl_kernel_launchI26rvalue_ref_arg_kernel_nameJZ4mainEUlT_E0_EEvPKcDpT0_(ptr noundef @.str.7, i32 %1)
+// CHECK-HOST-LINUX-NEXT: call void @_Z18sycl_kernel_launchI18__sycl_kernel_infoI26rvalue_ref_arg_kernel_nameEJZ4mainEUlT_E0_EEvDpT0_(i32 %1)
// CHECK-HOST-LINUX-NEXT: ret void
// CHECK-HOST-LINUX-NEXT: }
@@ -267,7 +250,7 @@ int main() {
// CHECK-HOST-WINDOWS-NEXT: store i8 %kernelFunc.coerce, ptr %coerce.dive, align 1
// 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: call void @"??$sycl_kernel_launch at V?$__sycl_kernel_info at Usingle_purpose_kernel_name@@@@Usingle_purpose_kernel@@@@YAXUsingle_purpose_kernel@@@Z"(i8 %0)
// CHECK-HOST-WINDOWS-NEXT: ret void
// CHECK-HOST-WINDOWS-NEXT: }
//
@@ -280,7 +263,7 @@ int main() {
// 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: call void @"??$sycl_kernel_launch at V?$__sycl_kernel_info at V<lambda_1>@?0??main@@9@@@V<lambda_1>@?0??main@@9@@@YAXV<lambda_1>@?0??main@@9@@Z"(i32 %0)
//
// CHECK-HOST-WINDOWS-NEXT: ret void
// CHECK-HOST-WINDOWS-NEXT: }
@@ -293,7 +276,7 @@ int main() {
// CHECK-HOST-WINDOWS-NEXT: store i8 %kernelFunc.coerce, ptr %coerce.dive, align 1
// 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: call void @"??$sycl_kernel_launch at V?$__sycl_kernel_info at U\CE\B4\CF\84\CF\87@@@@V<lambda_2>@?0??main@@9@@@YAXV<lambda_2>@?0??main@@9@@Z"(i8 %0)
// CHECK-HOST-WINDOWS-NEXT: ret void
// CHECK-HOST-WINDOWS-NEXT: }
@@ -317,7 +300,7 @@ int main() {
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive3 = getelementptr inbounds nuw %class.anon.1, ptr %agg.tmp, i32 0, i32 0
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive4 = getelementptr inbounds nuw %struct.copyable, ptr %coerce.dive3, i32 0, i32 0
// CHECK-HOST-WINDOWS-NEXT: %2 = load i32, ptr %coerce.dive4, align 4
-// CHECK-HOST-WINDOWS-NEXT: call void @"??$sycl_kernel_launch at UKN@?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) %this2, ptr noundef @"??_C at _0P@DLGHPODL at _ZTSZ4mainE2KN?$AA@", i32 %2, i32 noundef %1, i32 noundef %0)
+// CHECK-HOST-WINDOWS-NEXT: call void @"??$sycl_kernel_launch at V?$__sycl_kernel_info at UKN@?1??main@@9@@@V<lambda_3>@?0??main@@9 at HH@handler@@AEAAXV<lambda_3>@?0??main@@9 at HH@Z"(ptr noundef nonnull align 1 dereferenceable(1) %this2, i32 %2, i32 noundef %1, i32 noundef %0)
// CHECK-HOST-WINDOWS-NEXT: call void @"??1<lambda_3>@?0??main@@9 at QEAA@XZ"(ptr noundef nonnull align 4 dead_on_return(4) dereferenceable(4) %k) #{{[0-9]+}}
// CHECK-HOST-WINDOWS-NEXT: ret void
// CHECK-HOST-WINDOWS-NEXT: }
@@ -331,7 +314,7 @@ int main() {
// CHECK-HOST-WINDOWS-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 %agg.tmp, ptr align 4 %0, i64 4, i1 false)
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive = getelementptr inbounds nuw %class.anon, ptr %agg.tmp, i32 0, i32 0
// CHECK-HOST-WINDOWS-NEXT: %1 = load i32, ptr %coerce.dive, align 4
-// CHECK-HOST-WINDOWS-NEXT: call void @"??$sycl_kernel_launch at Uref_arg_kernel_name@@V<lambda_1>@?0??main@@9@@@YAXPEBDV<lambda_1>@?0??main@@9@@Z"(ptr noundef @"??_C at _0BK@PPDJPOBM at _ZTS19ref_arg_kernel_name?$AA@", i32 %1)
+// CHECK-HOST-WINDOWS-NEXT: call void @"??$sycl_kernel_launch at V?$__sycl_kernel_info at Uref_arg_kernel_name@@@@V<lambda_1>@?0??main@@9@@@YAXV<lambda_1>@?0??main@@9@@Z"(i32 %1)
// CHECK-HOST-WINDOWS-NEXT: ret void
// CHECK-HOST-WINDOWS-NEXT: }
//
@@ -344,7 +327,7 @@ int main() {
// CHECK-HOST-WINDOWS-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 %agg.tmp, ptr align 4 %0, i64 4, i1 false)
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive = getelementptr inbounds nuw %class.anon, ptr %agg.tmp, i32 0, i32 0
// CHECK-HOST-WINDOWS-NEXT: %1 = load i32, ptr %coerce.dive, align 4
-// CHECK-HOST-WINDOWS-NEXT: call void @"??$sycl_kernel_launch at Ufwd_ref_arg_kernel_name@@V<lambda_1>@?0??main@@9@@@YAXPEBDV<lambda_1>@?0??main@@9@@Z"(ptr noundef @"??_C at _0BO@KEIBIHKH at _ZTS23fwd_ref_arg_kernel_name?$AA@", i32 %1)
+// CHECK-HOST-WINDOWS-NEXT: call void @"??$sycl_kernel_launch at V?$__sycl_kernel_info at Ufwd_ref_arg_kernel_name@@@@V<lambda_1>@?0??main@@9@@@YAXV<lambda_1>@?0??main@@9@@Z"(i32 %1)
// CHECK-HOST-WINDOWS-NEXT: ret void
// CHECK-HOST-WINDOWS-NEXT: }
//
@@ -357,7 +340,7 @@ int main() {
// CHECK-HOST-WINDOWS-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 %agg.tmp, ptr align 4 %0, i64 4, i1 false)
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive = getelementptr inbounds nuw %class.anon, ptr %agg.tmp, i32 0, i32 0
// CHECK-HOST-WINDOWS-NEXT: %1 = load i32, ptr %coerce.dive, align 4
-// CHECK-HOST-WINDOWS-NEXT: call void @"??$sycl_kernel_launch at Ufwd_ref_arg_kernel_name_move@@V<lambda_1>@?0??main@@9@@@YAXPEBDV<lambda_1>@?0??main@@9@@Z"(ptr noundef @"??_C at _0CD@FDALJLMM at _ZTS28fwd_ref_arg_kernel_name_mo@", i32 %1)
+// CHECK-HOST-WINDOWS-NEXT: call void @"??$sycl_kernel_launch at V?$__sycl_kernel_info at Ufwd_ref_arg_kernel_name_move@@@@V<lambda_1>@?0??main@@9@@@YAXV<lambda_1>@?0??main@@9@@Z"(i32 %1)
// CHECK-HOST-WINDOWS-NEXT: ret void
// CHECK-HOST-WINDOWS-NEXT: }
//
@@ -370,7 +353,7 @@ int main() {
// CHECK-HOST-WINDOWS-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 4 %agg.tmp, ptr align 4 %0, i64 4, i1 false)
// CHECK-HOST-WINDOWS-NEXT: %coerce.dive = getelementptr inbounds nuw %class.anon.2, ptr %agg.tmp, i32 0, i32 0
// CHECK-HOST-WINDOWS-NEXT: %1 = load i32, ptr %coerce.dive, align 4
-// CHECK-HOST-WINDOWS-NEXT: call void @"??$sycl_kernel_launch at Urvalue_ref_arg_kernel_name@@V<lambda_4>@?0??main@@9@@@YAXPEBDV<lambda_4>@?0??main@@9@@Z"(ptr noundef @"??_C at _0CB@HCPMABHM at _ZTS26rvalue_ref_arg_kernel_name@", i32 %1)
+// CHECK-HOST-WINDOWS-NEXT: call void @"??$sycl_kernel_launch at V?$__sycl_kernel_info at Urvalue_ref_arg_kernel_name@@@@V<lambda_4>@?0??main@@9@@@YAXV<lambda_4>@?0??main@@9@@Z"(i32 %1)
// CHECK-HOST-WINDOWS-NEXT: ret void
// CHECK-HOST-WINDOWS-NEXT: }
diff --git a/clang/test/CodeGenSYCL/sycl-kernel-entry-point-exceptions.cpp b/clang/test/CodeGenSYCL/sycl-kernel-entry-point-exceptions.cpp
index 8fe7a148a2f61..11e19ad776fec 100644
--- a/clang/test/CodeGenSYCL/sycl-kernel-entry-point-exceptions.cpp
+++ b/clang/test/CodeGenSYCL/sycl-kernel-entry-point-exceptions.cpp
@@ -21,8 +21,8 @@ struct KT {
// calls a potentially throwing sycl_kernel_launch function (a thrown
// exception will propagate with no explicit handling required).
namespace ns1 {
- template<typename KN, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...);
+ template<typename KI, typename... Ts>
+ void sycl_kernel_launch(Ts...);
[[clang::sycl_kernel_entry_point(KN<1>)]]
void skep(KT<1> k) {
k();
@@ -30,7 +30,7 @@ namespace ns1 {
}
// CHECK: ; Function Attrs: mustprogress noinline optnone
// CHECK: define dso_local void @_ZN3ns14skepE2KTILi1ELi0EE() #{{[0-9]+}} {
-// CHECK: call void @_ZN3ns118sycl_kernel_launchI2KNILi1EEJ2KTILi1ELi0EEEEEvPKcDpT0_(ptr noundef @.str)
+// CHECK: call void @_ZN3ns118sycl_kernel_launchI18__sycl_kernel_infoI2KNILi1EEEJ2KTILi1ELi0EEEEEvDpT0_()
// CHECK: ret void
// CHECK: }
@@ -39,8 +39,8 @@ namespace ns1 {
// non-throwing sycl_kernel_entry_point attributed function calls
// a potentially throwing sycl_kernel_launch function.
namespace ns2 {
- template<typename KN, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...);
+ template<typename KI, typename... Ts>
+ void sycl_kernel_launch(Ts...);
[[clang::sycl_kernel_entry_point(KN<2>)]]
void skep(KT<2> k) noexcept {
k();
@@ -48,7 +48,7 @@ namespace ns2 {
}
// CHECK: ; Function Attrs: mustprogress noinline nounwind optnone
// CHECK: define dso_local void @_ZN3ns24skepE2KTILi2ELi0EE() #{{[0-9]+}} personality ptr @__gxx_personality_v0 {
-// CHECK: invoke void @_ZN3ns218sycl_kernel_launchI2KNILi2EEJ2KTILi2ELi0EEEEEvPKcDpT0_(ptr noundef @.str.1)
+// CHECK: invoke void @_ZN3ns218sycl_kernel_launchI18__sycl_kernel_infoI2KNILi2EEEJ2KTILi2ELi0EEEEEvDpT0_()
// CHECK: to label %invoke.cont unwind label %terminate.lpad
// CHECK: invoke.cont:
// CHECK: ret void
@@ -63,8 +63,8 @@ namespace ns2 {
// calls a non-throwing sycl_kernel_launch function (a thrown
// exception will terminate within sycl_kernel_launch).
namespace ns3 {
- template<typename KN, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...) noexcept;
+ template<typename KI, typename... Ts>
+ void sycl_kernel_launch(Ts...) noexcept;
[[clang::sycl_kernel_entry_point(KN<3>)]]
void skep(KT<3> k) {
k();
@@ -72,7 +72,7 @@ namespace ns3 {
}
// CHECK: ; Function Attrs: mustprogress noinline nounwind optnone
// CHECK: define dso_local void @_ZN3ns34skepE2KTILi3ELi0EE() #{{[0-9]+}} {
-// CHECK: call void @_ZN3ns318sycl_kernel_launchI2KNILi3EEJ2KTILi3ELi0EEEEEvPKcDpT0_(ptr noundef @.str.2)
+// CHECK: call void @_ZN3ns318sycl_kernel_launchI18__sycl_kernel_infoI2KNILi3EEEJ2KTILi3ELi0EEEEEvDpT0_()
// CHECK: ret void
// CHECK: }
@@ -81,8 +81,8 @@ namespace ns3 {
// non-throwing sycl_kernel_entry_point attributed function calls a
// non-throwing sycl_kernel_launch function.
namespace ns4 {
- template<typename KN, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...) noexcept;
+ template<typename KI, typename... Ts>
+ void sycl_kernel_launch(Ts...) noexcept;
[[clang::sycl_kernel_entry_point(KN<4>)]]
void skep(KT<4> k) noexcept {
k();
@@ -90,6 +90,6 @@ namespace ns4 {
}
// CHECK: ; Function Attrs: mustprogress noinline nounwind optnone
// CHECK: define dso_local void @_ZN3ns44skepE2KTILi4ELi0EE() #{{[0-9]+}} {
-// CHECK: call void @_ZN3ns418sycl_kernel_launchI2KNILi4EEJ2KTILi4ELi0EEEEEvPKcDpT0_(ptr noundef @.str.3)
+// CHECK: call void @_ZN3ns418sycl_kernel_launchI18__sycl_kernel_infoI2KNILi4EEEJ2KTILi4ELi0EEEEEvDpT0_()
// CHECK: ret void
// CHECK: }
diff --git a/clang/test/CodeGenSYCL/sycl-module-id.cpp b/clang/test/CodeGenSYCL/sycl-module-id.cpp
index d9beded116a65..890ee012cbe71 100644
--- a/clang/test/CodeGenSYCL/sycl-module-id.cpp
+++ b/clang/test/CodeGenSYCL/sycl-module-id.cpp
@@ -6,8 +6,8 @@
// 2. sycl_external functions
// Required by sycl_kernel_entry_point semantics.
-template <typename KernelName, typename... Ts>
-void sycl_kernel_launch(const char *, Ts...) {}
+template <typename KernelInfo, typename... Ts>
+void sycl_kernel_launch(Ts...) {}
template <typename KernelName, typename KernelType>
[[clang::sycl_kernel_entry_point(KernelName)]]
diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
index c298593e2f1ab..1862163aec81f 100644
--- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
+++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
@@ -1,8 +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 KN, typename... Ts>
-void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KI, typename... Ts>
+void sycl_kernel_launch(Ts...) {}
template<typename KN, typename Func>
[[clang::sycl_kernel_entry_point(KN)]] void kernel(Func F){
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 18118fc43a602..553a3fac128ce 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp
@@ -41,8 +41,8 @@ struct coroutine_traits {
template<int, int = 0> struct KN;
// A generic kernel launch function.
-template<typename KNT, typename... Ts>
-void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KIT, typename... Ts>
+void sycl_kernel_launch(Ts...) {}
////////////////////////////////////////////////////////////////////////////////
// Valid declarations.
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-device-odr-use.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-device-odr-use.cpp
index 631e1d7ad3870..8afc1b47476bc 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-device-odr-use.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-device-odr-use.cpp
@@ -17,8 +17,8 @@ struct type_info {
} // namespace std
// A generic kernel launch function.
-template<typename KernelName, typename... Ts>
-void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KernelInfo, typename... Ts>
+void sycl_kernel_launch(Ts...) {}
// A kernel name type template.
template<int> struct KN;
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 3a0c0d06daa5f..60d09cf0c2ab0 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
@@ -11,8 +11,8 @@ template<int> struct ST; // #ST-decl
template<int N> using TTA = ST<N>; // #TTA-decl
// A generic kernel launch function.
-template<typename KN, typename... Ts>
-void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KI, typename... Ts>
+void sycl_kernel_launch(Ts...) {}
////////////////////////////////////////////////////////////////////////////////
// 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 05a660e91e82c..f57764694e450 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
@@ -18,8 +18,8 @@ module M2 { header "m2.h" }
template<int> struct KN;
// A generic kernel launch function.
-template<typename KN, typename... Ts>
-void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KI, typename... Ts>
+void sycl_kernel_launch(Ts...) {}
[[clang::sycl_kernel_entry_point(KN<1>)]]
void common_test1() {}
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 dcea60e016d12..4216029d26e62 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
@@ -16,8 +16,8 @@
template<int> struct KN;
// A generic kernel launch function.
-template<typename KN, typename... Ts>
-void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KI, typename... Ts>
+void sycl_kernel_launch(Ts...) {}
[[clang::sycl_kernel_entry_point(KN<1>)]]
void pch_test1() {} // << expected previous declaration note here.
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 70d9e5f218139..c8aecb1a072c2 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
@@ -11,8 +11,8 @@
struct S1;
// A generic kernel launch function.
-template<typename KernelName, typename... Ts>
-void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KernelInfo, typename... Ts>
+void sycl_kernel_launch(Ts...) {}
// expected-warning at +3 {{redundant 'clang::sycl_kernel_entry_point' attribute}}
// expected-note at +1 {{previous attribute is here}}
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 654564f675e54..7cd7335be7ac4 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp
@@ -11,8 +11,8 @@
// is selected by overload resolution.
// A generic kernel launch function.
-template<typename KernelName, typename... Ts>
-void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KernelInfo, typename... Ts>
+void sycl_kernel_launch(Ts...) {}
// FIXME: C++23 [temp.expl.spec]p12 states:
// FIXME: ... Similarly, attributes appearing in the declaration of a 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 ac2d5c1541528..51a7ef6a7b9e0 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp
@@ -23,8 +23,8 @@ struct type_info {
} // namespace std
// A generic kernel launch function.
-template<typename KernelName, typename... Ts>
-void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KernelInfo, typename... Ts>
+void sycl_kernel_launch(Ts...) {}
////////////////////////////////////////////////////////////////////////////////
// Valid declarations.
diff --git a/clang/test/SemaSYCL/sycl-kernel-launch-ms-compat.cpp b/clang/test/SemaSYCL/sycl-kernel-launch-ms-compat.cpp
index f62b129273217..d607839bffc97 100644
--- a/clang/test/SemaSYCL/sycl-kernel-launch-ms-compat.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-launch-ms-compat.cpp
@@ -24,8 +24,8 @@ namespace ok1 {
struct base_handler {
protected:
// expected-note at +2 {{must qualify identifier to find this declaration in dependent base class}}
- template<typename KN, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...);
+ template<typename KI, typename... Ts>
+ void sycl_kernel_launch(Ts...);
};
template<int N>
struct handler : protected base_handler<handler<N>> {
@@ -34,7 +34,7 @@ namespace ok1 {
// lookups are allowed as a Microsoft compatible extension.
// expected-warning at +4 {{use of member 'sycl_kernel_launch' found via unqualified lookup into dependent bases of class templates is a Microsoft extension}}
// expected-note at +2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note-re at +1 {{in implicit call to 'sycl_kernel_launch' with template argument 'KN<1>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'KT<1>') required here}}
+ // expected-note at +1 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<KN<1>>' and function arguments (xvalue of type 'KT<1>') required here}}
[[clang::sycl_kernel_entry_point(KN<1>)]]
void skep(KT<1> k) {
k();
@@ -65,8 +65,8 @@ namespace bad1 {
private:
// expected-note at +3 {{must qualify identifier to find this declaration in dependent base class}}
// expected-note at +2 {{declared private here}}
- template<typename KN, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...);
+ template<typename KI, typename... Ts>
+ void sycl_kernel_launch(Ts...);
};
template<int N>
struct handler : protected base_handler<handler<N>> {
@@ -76,7 +76,7 @@ namespace bad1 {
// this case an error.
// expected-warning at +5 {{use of member 'sycl_kernel_launch' found via unqualified lookup into dependent bases of class templates is a Microsoft extension}}
// expected-note at +3 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note-re at +2 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<1>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<1>') required here}}
+ // expected-note at +2 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<BADKN<1>>' and function arguments (xvalue of type 'BADKT<1>') required here}}
// expected-error at +2 {{'sycl_kernel_launch' is a private member of 'bad1::base_handler<bad1::handler<1>>'}}
[[clang::sycl_kernel_entry_point(BADKN<1>)]]
void skep(BADKT<1> k) {
diff --git a/clang/test/SemaSYCL/sycl-kernel-launch.cpp b/clang/test/SemaSYCL/sycl-kernel-launch.cpp
index 06452de35b3d2..79d3ccc3b816f 100644
--- a/clang/test/SemaSYCL/sycl-kernel-launch.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-launch.cpp
@@ -5,7 +5,7 @@
// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++23 -fsyntax-only -fsycl-is-host -fcxx-exceptions -verify %s
// RUN: %clang_cc1 -triple spirv64-unknown-unknown -std=c++23 -fsyntax-only -fsycl-is-device -verify %s
-// Test overload resolution for implicit calls to sycl_kernel_launch<KN>(...)
+// Test overload resolution for implicit calls to sycl_kernel_launch<KI>(...)
// synthesized for functions declared with the sycl_kernel_entry_point
// attribute.
@@ -26,8 +26,8 @@ struct KT {
// sycl_kernel_launch as function template at namespace scope.
namespace ok1 {
- template<typename KN, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...);
+ template<typename KI, typename... Ts>
+ void sycl_kernel_launch(Ts...);
[[clang::sycl_kernel_entry_point(KN<1>)]]
void skep(KT<1> k) {
k();
@@ -37,8 +37,8 @@ namespace ok1 {
// sycl_kernel_launch as function template at namespace scope with default
// template arguments and default function arguments..
namespace ok2 {
- template<typename KN, typename T = int>
- void sycl_kernel_launch(const char *, KT<2>, T = 2);
+ template<typename KI, typename T = int>
+ void sycl_kernel_launch(KT<2>, T = 2);
[[clang::sycl_kernel_entry_point(KN<2>)]]
void skep(KT<2> k) {
k();
@@ -47,10 +47,10 @@ namespace ok2 {
// sycl_kernel_launch as overload set.
namespace ok3 {
- template<typename KN>
- void sycl_kernel_launch(const char *);
- template<typename KN, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...);
+ template<typename KI>
+ void sycl_kernel_launch(int);
+ template<typename KI, typename... Ts>
+ void sycl_kernel_launch(Ts...);
[[clang::sycl_kernel_entry_point(KN<3>)]]
void skep(KT<3> k) {
k();
@@ -61,8 +61,8 @@ namespace ok3 {
namespace ok4 {
struct handler {
private:
- template<typename KN, typename... Ts>
- static void sycl_kernel_launch(const char *, Ts...);
+ template<typename KI, typename... Ts>
+ static void sycl_kernel_launch(Ts...);
public:
[[clang::sycl_kernel_entry_point(KN<4,0>)]]
static void skep(KT<4,0> k) {
@@ -79,8 +79,8 @@ namespace ok4 {
namespace ok5 {
struct handler {
private:
- template<typename KN, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...);
+ template<typename KI, typename... Ts>
+ void sycl_kernel_launch(Ts...);
public:
[[clang::sycl_kernel_entry_point(KN<5>)]]
void skep(KT<5> k) {
@@ -95,8 +95,8 @@ namespace ok5 {
namespace ok6 {
struct handler {
private:
- template<typename KN, typename... Ts>
- void sycl_kernel_launch(this handler self, const char *, Ts...);
+ template<typename KI, typename... Ts>
+ void sycl_kernel_launch(this handler self, Ts...);
public:
[[clang::sycl_kernel_entry_point(KN<6>)]]
void skep(KT<6> k) {
@@ -108,13 +108,13 @@ namespace ok6 {
// sycl_kernel_launch as variable template.
namespace ok7 {
- template<typename KN>
+ template<typename KI>
struct launcher {
template<typename... Ts>
- void operator()(const char *, Ts...);
+ void operator()(Ts...);
};
- template<typename KN>
- launcher<KN> sycl_kernel_launch;
+ template<typename KI>
+ launcher<KI> sycl_kernel_launch;
[[clang::sycl_kernel_entry_point(KN<7>)]]
void skep(KT<7> k) {
k();
@@ -124,13 +124,13 @@ namespace ok7 {
#if __cplusplus >= 202302L
// sycl_kernel_launch as variable template with static call operator template.
namespace ok8 {
- template<typename KN>
+ template<typename KI>
struct launcher {
template<typename... Ts>
- static void operator()(const char *, Ts...);
+ static void operator()(Ts...);
};
- template<typename KN>
- launcher<KN> sycl_kernel_launch;
+ template<typename KI>
+ launcher<KI> sycl_kernel_launch;
[[clang::sycl_kernel_entry_point(KN<8>)]]
void skep(KT<8> k) {
k();
@@ -142,13 +142,13 @@ namespace ok8 {
// sycl_kernel_launch as variable template with call operator template with
// explicit object parameter.
namespace ok9 {
- template<typename KN>
+ template<typename KI>
struct launcher {
template<typename... Ts>
- void operator()(this launcher self, const char *, Ts...);
+ void operator()(this launcher self, Ts...);
};
- template<typename KN>
- launcher<KN> sycl_kernel_launch;
+ template<typename KI>
+ launcher<KI> sycl_kernel_launch;
[[clang::sycl_kernel_entry_point(KN<9>)]]
void skep(KT<9> k) {
k();
@@ -161,8 +161,8 @@ namespace ok10 {
template<typename Derived>
struct base_handler {
protected:
- template<typename KN, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...);
+ template<typename KI, typename... Ts>
+ void sycl_kernel_launch(Ts...);
};
struct handler : protected base_handler<handler> {
public:
@@ -175,8 +175,8 @@ namespace ok10 {
// sycl_kernel_launch with non-reference parameters.
namespace ok11 {
- template<typename KN, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...);
+ template<typename KI, typename... Ts>
+ void sycl_kernel_launch(Ts...);
struct move_only {
move_only(move_only&&) = default;
};
@@ -188,8 +188,8 @@ namespace ok11 {
// sycl_kernel_launch with forward reference parameters.
namespace ok12 {
- template<typename KN, typename... Ts>
- void sycl_kernel_launch(const char *, Ts &&...);
+ template<typename KI, typename... Ts>
+ void sycl_kernel_launch(Ts &&...);
struct non_copyable {
non_copyable(const non_copyable&) = delete;
};
@@ -213,8 +213,8 @@ namespace ok13 {
k();
}
namespace nested {
- template<typename KN, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...);
+ template<typename KI, typename... Ts>
+ void sycl_kernel_launch(Ts...);
struct S13 {};
}
template void skep<KN<13>>(KT<13>, nested::S13);
@@ -239,7 +239,7 @@ struct BADKT {
namespace bad1 {
// expected-error at +4 {{use of undeclared identifier 'sycl_kernel_launch'}}
// expected-note at +2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note-re at +1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<1>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<1>') required here}}
+ // expected-note at +1 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<BADKN<1>>' and function arguments (xvalue of type 'BADKT<1>') required here}}
[[clang::sycl_kernel_entry_point(BADKN<1>)]]
void skep(BADKT<1> k) {
k();
@@ -250,7 +250,7 @@ namespace bad1 {
namespace bad2 {
// expected-error at +5 {{use of undeclared identifier 'sycl_kernel_launch'}}
// expected-note at +3 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note-re at +2 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<2>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<2>') required here}}
+ // expected-note at +2 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<BADKN<2, 0>>' and function arguments (xvalue of type 'BADKT<2>') required here}}
template<typename KN, typename KT>
[[clang::sycl_kernel_entry_point(KN)]]
void skep(KT k) {
@@ -263,10 +263,10 @@ namespace bad2 {
// No matching function for call to sycl_kernel_launch; not a template.
namespace bad3 {
// expected-note at +1 {{declared as a non-template here}}
- void sycl_kernel_launch(const char *, BADKT<3>);
+ void sycl_kernel_launch(BADKT<3>);
// expected-error at +4 {{'sycl_kernel_launch' does not refer to a template}}
// expected-note at +2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note at +1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<3>' required here}}
+ // expected-note at +1 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<BADKN<3>>' required here}}
[[clang::sycl_kernel_entry_point(BADKN<3>)]]
void skep(BADKT<3> k) {
k();
@@ -275,12 +275,12 @@ namespace bad3 {
// No matching function for call to sycl_kernel_launch; not enough arguments.
namespace bad4 {
- // expected-note at +2 {{candidate function template not viable: requires 2 arguments, but 1 was provided}}
- template<typename KN, typename KT>
- void sycl_kernel_launch(const char *, KT);
+ // expected-note at +2 {{candidate function template not viable: requires 1 argument, but 0 were provided}}
+ template<typename KI>
+ void sycl_kernel_launch(int);
// expected-error at +5 {{no matching function for call to 'sycl_kernel_launch'}}
// expected-note at +3 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note-re at +2 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<4>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]') required here}}
+ // expected-note at +2 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<BADKN<4, 0>>' and function arguments () required here}}
template<typename KN>
[[clang::sycl_kernel_entry_point(KN)]]
void skep() {}
@@ -290,12 +290,12 @@ namespace bad4 {
// No matching function for call to sycl_kernel_launch; too many arguments.
namespace bad5 {
- // expected-note at +2 {{candidate function template not viable: requires 2 arguments, but 3 were provided}}
- template<typename KN, typename KT>
- void sycl_kernel_launch(const char *, KT);
+ // expected-note at +2 {{candidate function template not viable: requires 1 argument, but 2 were provided}}
+ template<typename KI, typename KT>
+ void sycl_kernel_launch(KT);
// expected-error at +5 {{no matching function for call to 'sycl_kernel_launch'}}
// expected-note at +3 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note-re at +2 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<5>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<5>', xvalue of type 'int') required here}}
+ // expected-note at +2 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<BADKN<5, 0>>' and function arguments (xvalue of type 'BADKT<5>', xvalue of type 'int') required here}}
template<typename KN, typename KT>
[[clang::sycl_kernel_entry_point(KN)]]
void skep(KT k, int i) {
@@ -307,12 +307,12 @@ namespace bad5 {
// No matching function for call to sycl_kernel_launch; mismatched function parameter type.
namespace bad6 {
- // expected-note-re at +2 {{candidate function template not viable: no known conversion from 'const char[{{[0-9]*}}]' to 'int' for 1st argument}}
- template<typename KN, typename... Ts>
+ // expected-note at +2 {{candidate function template not viable: no known conversion from 'BADKT<6>' to 'int' for 1st argument}}
+ template<typename KI, typename... Ts>
void sycl_kernel_launch(int, Ts...);
// expected-error at +5 {{no matching function for call to 'sycl_kernel_launch'}}
// expected-note at +3 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note-re at +2 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<6>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<6>') required here}}
+ // expected-note at +2 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<BADKN<6, 0>>' and function arguments (xvalue of type 'BADKT<6>') required here}}
template<typename KN, typename KT>
[[clang::sycl_kernel_entry_point(KN)]]
void skep(KT k) {
@@ -326,10 +326,10 @@ namespace bad6 {
namespace bad7 {
// expected-note at +2 {{candidate template ignored: template argument for non-type template parameter must be an expression}}
template<int, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...);
+ void sycl_kernel_launch(Ts...);
// expected-error at +4 {{no matching function for call to 'sycl_kernel_launch'}}
// expected-note at +2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note-re at +1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<7>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<7>') required here}}
+ // expected-note at +1 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<BADKN<7>>' and function arguments (xvalue of type 'BADKT<7>') required here}}
[[clang::sycl_kernel_entry_point(BADKN<7>)]]
void skep(BADKT<7> k) {
k();
@@ -338,12 +338,12 @@ namespace bad7 {
// No matching function for call to sycl_kernel_launch; substitution failure.
namespace bad8 {
- // expected-note at +2 {{candidate template ignored: substitution failure [with KN = BADKN<8>, KT = BADKT<8>]: no type named 'no_such_type' in 'BADKT<8>'}}
- template<typename KN, typename KT, typename T = typename KT::no_such_type>
- void sycl_kernel_launch(const char *, KT);
+ // expected-note at +2 {{candidate template ignored: substitution failure [with KI = __sycl_kernel_info<BADKN<8, 0>>, KT = BADKT<8>]: no type named 'no_such_type' in 'BADKT<8>'}}
+ template<typename KI, typename KT, typename T = typename KT::no_such_type>
+ void sycl_kernel_launch(KT);
// expected-error at +4 {{no matching function for call to 'sycl_kernel_launch'}}
// expected-note at +2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note-re at +1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<8>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<8>') required here}}
+ // expected-note at +1 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<BADKN<8>>' and function arguments (xvalue of type 'BADKT<8>') required here}}
[[clang::sycl_kernel_entry_point(BADKN<8>)]]
void skep(BADKT<8> k) {
k();
@@ -353,11 +353,11 @@ namespace bad8 {
// No matching function for call to sycl_kernel_launch; deduction failure.
namespace bad9 {
// expected-note at +2 {{candidate template ignored: couldn't infer template argument 'T'}}
- template<typename KN, typename KT, typename T>
- void sycl_kernel_launch(const char *, KT);
+ template<typename KI, typename KT, typename T>
+ void sycl_kernel_launch(KT);
// expected-error at +4 {{no matching function for call to 'sycl_kernel_launch'}}
// expected-note at +2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note-re at +1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<9>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<9>') required here}}
+ // expected-note at +1 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<BADKN<9>>' and function arguments (xvalue of type 'BADKT<9>') required here}}
[[clang::sycl_kernel_entry_point(BADKN<9>)]]
void skep(BADKT<9> k) {
k();
@@ -366,17 +366,17 @@ namespace bad9 {
// No matching function for call to sycl_kernel_launch object; mismatched function parameter type.
namespace bad10 {
- template<typename KN>
+ template<typename KI>
struct launcher {
- // expected-note-re at +2 {{candidate function template not viable: no known conversion from 'const char[{{[0-9]*}}]' to 'int' for 1st argument}}
+ // expected-note at +2 {{candidate function template not viable: no known conversion from 'BADKT<10>' to 'int' for 1st argument}}
template<typename... Ts>
void operator()(int, Ts...);
};
- template<typename KN>
- launcher<KN> sycl_kernel_launch;
- // expected-error at +5 {{no matching function for call to object of type 'launcher<BADKN<10, 0>>'}}
+ template<typename KI>
+ launcher<KI> sycl_kernel_launch;
+ // expected-error at +5 {{no matching function for call to object of type 'launcher<__sycl_kernel_info<BADKN<10, 0>>>'}}
// expected-note at +3 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note-re at +2 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<10>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<10>') required here}}
+ // expected-note at +2 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<BADKN<10, 0>>' and function arguments (xvalue of type 'BADKT<10>') required here}}
template<typename KN, typename KT>
[[clang::sycl_kernel_entry_point(KN)]]
void skep(KT k) {
@@ -388,17 +388,17 @@ namespace bad10 {
// No matching function for call to sycl_kernel_launch object; mismatched template parameter kind.
namespace bad11 {
- template<int KN>
+ template<int KI>
struct launcher {
template<typename... Ts>
- void operator()(int, Ts...);
+ void operator()(Ts...);
};
// expected-note at +1 {{template parameter is declared here}}
- template<int KN>
- launcher<KN> sycl_kernel_launch;
+ template<int KI>
+ launcher<KI> sycl_kernel_launch;
// expected-error at +5 {{template argument for non-type template parameter must be an expression}}
// expected-note at +3 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note at +2 {{in implicit call to 'sycl_kernel_launch' with template argument 'KN' required here}}
+ // expected-note at +2 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<KN>' required here}}
template<typename KN, typename KT>
[[clang::sycl_kernel_entry_point(KN)]]
void skep(KT k) {
@@ -409,18 +409,18 @@ namespace bad11 {
// sycl_kernel_launch as variable template with private call operator template.
namespace bad12 {
- template<typename KN>
+ template<typename KI>
struct launcher {
private:
// expected-note at +2 {{declared private here}}
template<typename... Ts>
- void operator()(const char *, Ts...);
+ void operator()(Ts...);
};
- template<typename KN>
- launcher<KN> sycl_kernel_launch;
- // expected-error at +4 {{'operator()' is a private member of 'bad12::launcher<BADKN<12>>'}}
+ template<typename KI>
+ launcher<KI> sycl_kernel_launch;
+ // expected-error at +4 {{'operator()' is a private member of 'bad12::launcher<__sycl_kernel_info<BADKN<12>>>'}}
// expected-note at +2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note-re at +1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<12>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<12>') required here}}
+ // expected-note at +1 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<BADKN<12>>' and function arguments (xvalue of type 'BADKT<12>') required here}}
[[clang::sycl_kernel_entry_point(BADKN<12>)]]
void skep(BADKT<12> k) {
k();
@@ -431,22 +431,22 @@ namespace bad12 {
namespace bad13 {
inline namespace in1 {
// expected-note at +2 {{candidate found by name lookup is 'bad13::in1::sycl_kernel_launch'}}
- template<typename KN, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...);
+ template<typename KI, typename... Ts>
+ void sycl_kernel_launch(Ts...);
}
inline namespace in2 {
- template<typename KN>
+ template<typename KI>
struct launcher {
template<typename KT, typename... Ts>
- void operator()(const char *, Ts...);
+ void operator()(Ts...);
};
// expected-note at +2 {{candidate found by name lookup is 'bad13::in2::sycl_kernel_launch'}}
- template<typename KN>
- launcher<KN> sycl_kernel_launch;
+ template<typename KI>
+ launcher<KI> sycl_kernel_launch;
}
// expected-error at +5 {{reference to 'sycl_kernel_launch' is ambiguous}}
// expected-note at +3 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note at +2 {{in implicit call to 'sycl_kernel_launch' with template argument 'KN' required here}}
+ // expected-note at +2 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<KN>' required here}}
template<typename KN, typename KT>
[[clang::sycl_kernel_entry_point(KN)]]
void skep(KT k) {
@@ -457,15 +457,15 @@ namespace bad13 {
// Ambiguous call to sycl_kernel_launch.
namespace bad14 {
- // expected-note at +2 {{candidate function [with KN = BADKN<14>, KT = BADKT<14>]}}
- template<typename KN, typename KT>
- void sycl_kernel_launch(const char *, KT, signed char);
- // expected-note at +2 {{candidate function [with KN = BADKN<14>, KT = BADKT<14>]}}
- template<typename KN, typename KT>
- void sycl_kernel_launch(const char *, KT, unsigned char);
+ // expected-note at +2 {{candidate function [with KI = __sycl_kernel_info<BADKN<14>>, KT = BADKT<14>]}}
+ template<typename KI, typename KT>
+ void sycl_kernel_launch(KT, signed char);
+ // expected-note at +2 {{candidate function [with KI = __sycl_kernel_info<BADKN<14>>, KT = BADKT<14>]}}
+ template<typename KI, typename KT>
+ void sycl_kernel_launch(KT, unsigned char);
// expected-error at +4 {{call to 'sycl_kernel_launch' is ambiguous}}
// expected-note at +2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note-re at +1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<14>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<14>', xvalue of type 'int') required here}}
+ // expected-note at +1 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<BADKN<14>>' and function arguments (xvalue of type 'BADKT<14>', xvalue of type 'int') required here}}
[[clang::sycl_kernel_entry_point(BADKN<14>)]]
void skep(BADKT<14> k, int i) {
k();
@@ -475,11 +475,11 @@ namespace bad14 {
// Call to member sycl_kernel_launch from non-static member.
namespace bad15 {
struct S {
- template<typename KN, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...);
+ template<typename KI, typename... Ts>
+ void sycl_kernel_launch(Ts...);
// expected-error at +4 {{call to non-static member function without an object argument}}
// expected-note at +2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note at +1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<15>' required here}}
+ // expected-note at +1 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<BADKN<15>>' required here}}
[[clang::sycl_kernel_entry_point(BADKN<15>)]]
static void skep(BADKT<15> k) {
k();
@@ -494,8 +494,8 @@ namespace bad16 {
struct base_handler {
protected:
// expected-note at +2 {{member is declared here}}
- template<typename KN, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...);
+ template<typename KI, typename... Ts>
+ void sycl_kernel_launch(Ts...);
};
template<int N>
struct handler : protected base_handler<handler<N>> {
@@ -503,7 +503,7 @@ namespace bad16 {
// classes requires explicit qualification.
// expected-error at +4 {{explicit qualification required to use member 'sycl_kernel_launch' from dependent base class}}
// expected-note at +2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note-re at +1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<16>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<16>') required here}}
+ // expected-note at +1 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<BADKN<16>>' and function arguments (xvalue of type 'BADKT<16>') required here}}
[[clang::sycl_kernel_entry_point(BADKN<16>)]]
void skep(BADKT<16> k) {
k();
@@ -516,15 +516,15 @@ namespace bad16 {
// sycl_kernel_launch with non-reference parameters and non-moveable arguments.
namespace bad17 {
// expected-note at +2 2 {{passing argument to parameter here}}
- template<typename KN, typename... Ts>
- void sycl_kernel_launch(const char *, Ts...);
+ template<typename KI, typename... Ts>
+ void sycl_kernel_launch(Ts...);
struct non_copyable {
// expected-note at +1 {{'non_copyable' has been explicitly marked deleted here}}
non_copyable(const non_copyable&) = delete;
};
// expected-error at +4 {{call to deleted constructor of 'bad17::non_copyable'}}
// expected-note at +2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note-re at +1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<17, 0>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<17, 0>', xvalue of type 'non_copyable') required here}}
+ // expected-note at +1 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<BADKN<17, 0>>' and function arguments (xvalue of type 'BADKT<17, 0>', xvalue of type 'non_copyable') required here}}
[[clang::sycl_kernel_entry_point(BADKN<17,0>)]]
void skep(BADKT<17,0> k, non_copyable) {
k();
@@ -535,7 +535,7 @@ namespace bad17 {
};
// expected-error at +4 {{call to deleted constructor of 'bad17::non_moveable'}}
// expected-note at +2 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note-re at +1 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<17, 1>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<17, 1>', xvalue of type 'non_moveable') required here}}
+ // expected-note at +1 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<BADKN<17, 1>>' and function arguments (xvalue of type 'BADKT<17, 1>', xvalue of type 'non_moveable') required here}}
[[clang::sycl_kernel_entry_point(BADKN<17,1>)]]
void skep(BADKT<17,1> k, non_moveable) {
k();
@@ -546,14 +546,14 @@ namespace bad17 {
namespace bad18 {
// expected-error at +5 {{call to function 'sycl_kernel_launch' that is neither visible in the template definition nor found by argument-dependent lookup}}
// expected-note at +3 {{this indicates a problem with the SYCL runtime header files; please consider reporting this to your SYCL runtime provider}}
- // expected-note-re at +2 {{in implicit call to 'sycl_kernel_launch' with template argument 'BADKN<18>' and function arguments (lvalue of type 'const char[{{[0-9]*}}]', xvalue of type 'BADKT<18>') required here}}
+ // expected-note at +2 {{in implicit call to 'sycl_kernel_launch' with template argument '__sycl_kernel_info<BADKN<18, 0>>' and function arguments (xvalue of type 'BADKT<18>') required here}}
template<typename KN, typename KT>
[[clang::sycl_kernel_entry_point(KN)]]
void skep(KT k) {
k();
}
// expected-note at +2 {{'sycl_kernel_launch' should be declared prior to the call site or in the global namespace}}
- template<typename KN, typename... Ts>
+ template<typename KI, typename... Ts>
void sycl_kernel_launch(Ts...) {}
// expected-note at +1 {{in instantiation of function template specialization 'bad18::skep<BADKN<18>, BADKT<18>>' requested here}}
template void skep<BADKN<18>>(BADKT<18>);
diff --git a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
index 66aa00da18a04..6431f60773b66 100644
--- a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
@@ -5,8 +5,8 @@
template<int, int = 0> struct KN;
// A generic kernel launch function.
-template<typename KNT, typename... Ts>
-void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KIT, typename... Ts>
+void sycl_kernel_launch(Ts...) {}
// Check that reference captures of kernel that defined as lambda are diagnosed.
namespace badref1 {
diff --git a/libsycl/include/sycl/__impl/queue.hpp b/libsycl/include/sycl/__impl/queue.hpp
index 96d764e6f184a..d261a51f61e53 100644
--- a/libsycl/include/sycl/__impl/queue.hpp
+++ b/libsycl/include/sycl/__impl/queue.hpp
@@ -401,18 +401,17 @@ class _LIBSYCL_EXPORT queue {
/// Name of this function is defined by compiler. It generates a call to this
/// function in the host implementation of KernelFunc in submitSingleTask or
/// submitParallelFor.
- /// \param KernelName the name of the kernel being invoked.
/// \param args the kernel arguments for the kernel invocation.
- template <typename KN, typename... Args>
- void sycl_kernel_launch(const char *KernelName, Args &&...args) {
- static_assert(
- sizeof...(args) == 1,
- "sycl_kernel_launch expects only 2 arguments now: name of kernel and "
- "callable object passed to kernel invocation by the user.");
+ template <typename KI, typename... Args>
+ void sycl_kernel_launch(Args &&...args) {
+ static_assert(sizeof...(args) == 1,
+ "sycl_kernel_launch expects only 1 argument now: the "
+ "callable object passed to kernel invocation by the user.");
auto FirstArg = std::get<0>(std::tie(args...));
- submitKernelImpl(detail::getDeviceKernelInfo<KN>(KernelName), &FirstArg,
- sizeof(FirstArg));
+ submitKernelImpl(detail::getDeviceKernelInfo<typename KI::kernel_name>(
+ KI::kernel_entry_point_name),
+ &FirstArg, sizeof(FirstArg));
}
/// The sycl_kernel_entry_point attribute facilitates the generation of an
More information about the cfe-commits
mailing list