[clang] [clang][SYCL] Add support for handling of special SYCL types (PR #210319)
Mariya Podchishchaeva via cfe-commits
cfe-commits at lists.llvm.org
Fri Jul 17 05:55:08 PDT 2026
https://github.com/Fznamznon created https://github.com/llvm/llvm-project/pull/210319
Some SYCL types cannot be simply copied from host to device and instead require additional handling. Examples of such types are accessor/local_accessor/sampler/stream and etc.
This PR implements an interface provided by clang frontend to SYCL library implementers for special types handling. Previously implemented kernel launching mechanism implemented by https://github.com/llvm/llvm-project/pull/152403 is now extended with possibility to communicate which types within kernel parameters require special handling and implement it.
First, an attribute `clang::sycl_special_kernel_parameter` is introduced for marking types that require special handling within SYCL library. For example:
```
class [[clang::sycl_special_kernel_parameter]] accessor {
// code ...
};
```
Second, the result of the ``sycl_kernel_launch`` call now can either void or a function object (e.g., the result of a lambda expression, potentially one that captures references to the kernel arguments) for special kernel parameters processing purposes.
The special subobjects are then passed as lvalues in an invocation of the resulting function object.
Consider the following SYCL example
```
struct KN;
void f(sycl::handler &h, sycl::stream &sout, int i) {
h.single_task<KN>([=] {
sout << "The value of i is " << i << "\n";
});
}
```
A possible implementation of handler class could look like
```
namespace sycl {
class handler {
template <typename KernelName, typename KernelType>
[[clang::sycl_kernel_entry_point(KernelName)]]
void kernel_entry_point(KernelType kernelFunc) {
kernelFunc();
}
public:
template <typename KernelName, typename KernelType>
void single_task(const KernelType &kernelFunc) {
kernel_entry_point<KernelName>(kernelFunc);
}
};
}
```
For the example above, the following host pseudo code will be synthesized on AST level after this PR and will replace body of the function marked with `[[clang::sycl_kernel_entry_point(KernelName)]]` attribute:
```
sycl_kernel_launch<KernelName>("kernel-entry-point", kernelFunc)(kernelFunc.sout)
```
that generated code suffices to inform SYCL library about special kernel parameters, the task of SYCL library is to provide additional kernel parameter types required for initialization of mentioned special kernel parameters. This is done via returning an object of type `detail::type_list<buffer_t*, int>` from the call to the callable object returned by `sycl_kernel_launch`. Clang is now informed of two additional kernel parameters of types `buffer_t*` and `int` and can generate kernel signature.
Third, and last remaining part this PR brings is generation of actual processing of additional kernel parameters on the device side. This is done in a similar way to `sycl_kernel_launch` call generation via delegation to library-declared function template named `sycl_handle_special_kernel_parameters`. It is implicitly called at the beginning of the synthesized kernel and accepts references to special kernel parameters as direct arguments, returns a callable object which is immediately called with additional special kernel arguments (`buffer_t*, int` from the earlier example). The following pseudo code illustrates synthesized kernel
```
void kernel-entry-point(lambda-from-f kernelFunc, buffer_t* X, int Y) {
sycl_handle_special_kernel_parameters<kernel-name-type>(kernelFunc.sout)(X, Y);
kernelFunc();
}
```
Please refer to https://github.com/llvm/llvm-project/pull/170602 for additional details and explanations of the design.
>From 729266e9b5c7f0a39b5bcf31f52d03fce835f1f7 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Tue, 24 Feb 2026 14:03:17 -0800
Subject: [PATCH 01/11] [clang][SYCL] Decomposition support
WIP
---
clang/include/clang/AST/ASTNodeTraverser.h | 4 +-
clang/include/clang/AST/RecursiveASTVisitor.h | 1 +
clang/include/clang/AST/StmtSYCL.h | 26 +-
clang/include/clang/Basic/Attr.td | 7 +
clang/include/clang/Sema/ScopeInfo.h | 4 +
clang/include/clang/Sema/SemaSYCL.h | 9 +-
clang/lib/Sema/SemaDecl.cpp | 18 +-
clang/lib/Sema/SemaDeclAttr.cpp | 3 +
clang/lib/Sema/SemaSYCL.cpp | 297 ++++++++++++++++--
clang/lib/Sema/TreeTransform.h | 7 +-
.../ast-dump-sycl-kernel-call-stmt.cpp | 3 +
.../ast-dump-sycl-kernel-decomposition.cpp | 141 +++++++++
.../ast-dump-sycl-kernel-entry-point.cpp | 3 +
.../ASTSYCL/ast-print-sycl-kernel-call.cpp | 3 +
clang/test/CodeGenSYCL/function-attrs.cpp | 3 +
.../CodeGenSYCL/kernel-arg-decomposition.cpp | 96 ++++++
.../CodeGenSYCL/kernel-caller-entry-point.cpp | 3 +
.../sycl-kernel-entry-point-exceptions.cpp | 3 +
.../unique_stable_name_windows_diff.cpp | 3 +
...-kernel-entry-point-attr-appertainment.cpp | 3 +
...kernel-entry-point-attr-device-odr-use.cpp | 3 +
.../sycl-kernel-entry-point-attr-grammar.cpp | 3 +
...el-entry-point-attr-kernel-name-module.cpp | 3 +
...ernel-entry-point-attr-kernel-name-pch.cpp | 3 +
...cl-kernel-entry-point-attr-kernel-name.cpp | 3 +
.../sycl-kernel-entry-point-attr-sfinae.cpp | 3 +
.../sycl-kernel-entry-point-attr-this.cpp | 3 +
.../SemaSYCL/sycl-kernel-launch-ms-compat.cpp | 3 +
clang/test/SemaSYCL/sycl-kernel-launch.cpp | 3 +
29 files changed, 616 insertions(+), 48 deletions(-)
create mode 100644 clang/test/ASTSYCL/ast-dump-sycl-kernel-decomposition.cpp
create mode 100644 clang/test/CodeGenSYCL/kernel-arg-decomposition.cpp
diff --git a/clang/include/clang/AST/ASTNodeTraverser.h b/clang/include/clang/AST/ASTNodeTraverser.h
index d184e03355077..7c2d2413a4888 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->getKernelLaunchIdExpr());
+ Visit(Node->getSpecArgsIdExpr());
+ }
}
void VisitOMPExecutableDirective(const OMPExecutableDirective *Node) {
diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h
index b77443f1fa0ab..bc3ea412da756 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -3030,6 +3030,7 @@ DEF_TRAVERSE_STMT(UnresolvedSYCLKernelCallStmt, {
if (getDerived().shouldVisitImplicitCode()) {
TRY_TO(TraverseStmt(S->getOriginalStmt()));
TRY_TO(TraverseStmt(S->getKernelLaunchIdExpr()));
+ TRY_TO(TraverseStmt(S->getSpecArgsIdExpr()));
ShouldVisitChildren = false;
}
})
diff --git a/clang/include/clang/AST/StmtSYCL.h b/clang/include/clang/AST/StmtSYCL.h
index 79ac88532e143..cd682f4cea594 100644
--- a/clang/include/clang/AST/StmtSYCL.h
+++ b/clang/include/clang/AST/StmtSYCL.h
@@ -105,12 +105,19 @@ class UnresolvedSYCLKernelCallStmt : public Stmt {
Stmt *OriginalStmt = nullptr;
// KernelLaunchIdExpr stores an UnresolvedLookupExpr or UnresolvedMemberExpr
// corresponding to the SYCL kernel launch function for which a call
- // will be synthesized during template instantiation.
+ // will be synthesized during template instantiation of the host code.
Expr *KernelLaunchIdExpr = nullptr;
-
- UnresolvedSYCLKernelCallStmt(CompoundStmt *CS, Expr *IdExpr)
+ // Similar to KernelLaunchIdExpr HandleSYCLSpecialParamsIdExpr stores an
+ // UnresolvedLookupExpr or UnresolvedMemberExpr corresponding to the fuction
+ // handling of special SYCL kernel parameters for which a call will be
+ // synthesized during template instantiation of the device code.
+ Expr *HandleSYCLSpecialParamsIdExpr = nullptr;
+
+ UnresolvedSYCLKernelCallStmt(CompoundStmt *CS, Expr *IdExpr,
+ Expr *HandleSYCLSpecialParamsIdExpr)
: Stmt(UnresolvedSYCLKernelCallStmtClass), OriginalStmt(CS),
- KernelLaunchIdExpr(IdExpr) {}
+ KernelLaunchIdExpr(IdExpr),
+ HandleSYCLSpecialParamsIdExpr(HandleSYCLSpecialParamsIdExpr) {}
void setOriginalStmt(CompoundStmt *CS) { OriginalStmt = CS; }
@@ -118,12 +125,13 @@ class UnresolvedSYCLKernelCallStmt : public Stmt {
public:
static UnresolvedSYCLKernelCallStmt *Create(const ASTContext &C,
- CompoundStmt *CS, Expr *IdExpr) {
- return new (C) UnresolvedSYCLKernelCallStmt(CS, IdExpr);
+ CompoundStmt *CS, Expr *IdExpr,
+ Expr *SpecArgsExpr) {
+ return new (C) UnresolvedSYCLKernelCallStmt(CS, IdExpr, SpecArgsExpr);
}
static UnresolvedSYCLKernelCallStmt *CreateEmpty(const ASTContext &C) {
- return new (C) UnresolvedSYCLKernelCallStmt(nullptr, nullptr);
+ return new (C) UnresolvedSYCLKernelCallStmt(nullptr, nullptr, nullptr);
}
CompoundStmt *getOriginalStmt() { return cast<CompoundStmt>(OriginalStmt); }
@@ -133,6 +141,10 @@ class UnresolvedSYCLKernelCallStmt : public Stmt {
Expr *getKernelLaunchIdExpr() { return KernelLaunchIdExpr; }
const Expr *getKernelLaunchIdExpr() const { return KernelLaunchIdExpr; }
+ Expr *getSpecArgsIdExpr() { return HandleSYCLSpecialParamsIdExpr; }
+ const Expr *getSpecArgsIdExpr() const {
+ return HandleSYCLSpecialParamsIdExpr;
+ }
SourceLocation getBeginLoc() const LLVM_READONLY {
return getOriginalStmt()->getBeginLoc();
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 3f57104d474a7..7eb04ee1ffd9f 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1747,6 +1747,13 @@ def SYCLSpecialClass: InheritableAttr {
let Documentation = [SYCLSpecialClassDocs];
}
+def SYCLSpecialKernelParameter : InheritableAttr {
+ let Spellings = [CXX11<"clang", "sycl_special_kernel_parameter">];
+ let Subjects = SubjectList<[CXXRecord]>;
+ let LangOpts = [SYCLHost, SYCLDevice];
+ let Documentation = [Undocumented];
+}
+
def C11NoReturn : InheritableAttr {
let Spellings = [CustomKeyword<"_Noreturn">];
let Subjects = SubjectList<[Function], ErrorDiag>;
diff --git a/clang/include/clang/Sema/ScopeInfo.h b/clang/include/clang/Sema/ScopeInfo.h
index 7e4d3f2e0d1cb..6f9c0bf4c068c 100644
--- a/clang/include/clang/Sema/ScopeInfo.h
+++ b/clang/include/clang/Sema/ScopeInfo.h
@@ -249,6 +249,10 @@ class FunctionScopeInfo {
/// to a SYCL kernel launch function in a dependent context.
Expr *SYCLKernelLaunchIdExpr = nullptr;
+ /// An unresolved identifier lookup expression for an implicit call
+ /// to a handling function for SYCL kernel special parameters.
+ Expr *HandleSYCLSpecialParamsIdExpr = nullptr;
+
public:
/// Represents a simple identification of a weak object.
///
diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h
index 4980aa44c3012..268f31d8947cb 100644
--- a/clang/include/clang/Sema/SemaSYCL.h
+++ b/clang/include/clang/Sema/SemaSYCL.h
@@ -83,19 +83,22 @@ class SemaSYCL : public SemaBase {
/// 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);
+ ExprResult BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD, QualType KernelName,
+ StringRef FuncName);
/// 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);
+ Expr *LaunchIdExpr,
+ Expr *HandleSpecParamsExpr);
/// Builds an UnresolvedSYCLKernelCallStmt to wrap 'Body'. 'LaunchIdExpr'
/// specifies the lookup result returned by a previous call to
/// BuildSYCLKernelLaunchIdExpr().
StmtResult BuildUnresolvedSYCLKernelCallStmt(CompoundStmt *Body,
- Expr *LaunchIdExpr);
+ Expr *LaunchIdExpr,
+ Expr *HandleSpecParamsExpr);
};
} // namespace clang
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index f6d1d7a7877c7..48ff8eb7fc81f 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -16488,7 +16488,7 @@ Decl *Sema::ActOnStartOfFunctionDef(Scope *FnBodyScope, Decl *D,
const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>();
if (!SKEPAttr->isInvalidAttr()) {
ExprResult LaunchIdExpr =
- SYCL().BuildSYCLKernelLaunchIdExpr(FD, SKEPAttr->getKernelName());
+ SYCL().BuildSYCLKernelLaunchIdExpr(FD, SKEPAttr->getKernelName(), "sycl_kernel_launch");
// 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
@@ -16497,6 +16497,13 @@ Decl *Sema::ActOnStartOfFunctionDef(Scope *FnBodyScope, Decl *D,
// 'LaunchIDExpr' failed, then 'SYCLKernelLaunchIdExpr' will be assigned
// a null pointer value below; that is expected.
getCurFunction()->SYCLKernelLaunchIdExpr = LaunchIdExpr.get();
+ if (!LaunchIdExpr.isInvalid() &&
+ !LaunchIdExpr.get()->getType()->isVoidType()) {
+ ExprResult HSPSPIdExpr = SYCL().BuildSYCLKernelLaunchIdExpr(
+ FD, SKEPAttr->getKernelName(),
+ "sycl_handle_special_kernel_parameters");
+ getCurFunction()->HandleSYCLSpecialParamsIdExpr = HSPSPIdExpr.get();
+ }
}
}
@@ -16712,7 +16719,8 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body, bool IsInstantiation,
// 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) {
+ } else if (!getCurFunction()->SYCLKernelLaunchIdExpr ||
+ !getCurFunction()->HandleSYCLSpecialParamsIdExpr) {
// 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
@@ -16720,11 +16728,13 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body, bool IsInstantiation,
SR = Body;
} else if (FD->isTemplated()) {
SR = SYCL().BuildUnresolvedSYCLKernelCallStmt(
- cast<CompoundStmt>(Body), getCurFunction()->SYCLKernelLaunchIdExpr);
+ cast<CompoundStmt>(Body), getCurFunction()->SYCLKernelLaunchIdExpr,
+ getCurFunction()->HandleSYCLSpecialParamsIdExpr);
} else {
SR = SYCL().BuildSYCLKernelCallStmt(
FD, cast<CompoundStmt>(Body),
- getCurFunction()->SYCLKernelLaunchIdExpr);
+ getCurFunction()->SYCLKernelLaunchIdExpr,
+ getCurFunction()->HandleSYCLSpecialParamsIdExpr);
}
// 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/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 2159c586e5738..53f9dd0d49379 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -7774,6 +7774,9 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL,
case ParsedAttr::AT_SYCLSpecialClass:
handleSimpleAttribute<SYCLSpecialClassAttr>(S, D, AL);
break;
+ case ParsedAttr::AT_SYCLSpecialKernelParameter:
+ handleSimpleAttribute<SYCLSpecialKernelParameterAttr>(S, D, AL);
+ break;
case ParsedAttr::AT_Format:
handleFormatAttr(S, D, AL);
break;
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index b942f19761f40..62d9aa653300b 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -426,7 +426,8 @@ void SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD) {
}
ExprResult SemaSYCL::BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD,
- QualType KNT) {
+ QualType KNT,
+ StringRef FuncName) {
// 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 "
@@ -441,12 +442,13 @@ ExprResult SemaSYCL::BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD,
ASTContext &Ctx = SemaRef.getASTContext();
IdentifierInfo &SYCLKernelLaunchID =
- Ctx.Idents.get("sycl_kernel_launch", tok::TokenKind::identifier);
+ Ctx.Idents.get(FuncName, tok::TokenKind::identifier);
// Establish a code synthesis context for the implicit name lookup of
// a template named 'sycl_kernel_launch'. In the event of an error, this
// ensures an appropriate diagnostic note is issued to explain why the
// lookup was performed.
+ // FIXME: Extend diagnostics for handle special parameters function
Sema::CodeSynthesisContext CSC;
CSC.Kind = Sema::CodeSynthesisContext::SYCLKernelLaunchLookup;
CSC.Entity = FD;
@@ -493,16 +495,130 @@ ExprResult SemaSYCL::BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD,
return IdExpr;
}
+static bool isSyclSpecialType(QualType Ty) {
+ if (const auto *RT = Ty->getAsRecordDecl())
+ return RT->getMostRecentDecl()->hasAttr<SYCLSpecialKernelParameterAttr>();
+ return false;
+}
+
namespace {
+class KernelSpecialParamsCreator
+ : public SubobjectVisitor<KernelSpecialParamsCreator> {
+ using ObjectAccess =
+ llvm::PointerUnion<ValueDecl *, CXXBaseSpecifier *, FieldDecl *>;
+ SemaSYCL &SemaSYCLRef;
+ SourceLocation SrcLoc;
+ SmallVector<ObjectAccess, 4> ObjectAccessPath;
+ llvm::SmallVectorImpl<Expr *> &ResultingArgs;
+ RecordDecl *CurRD;
+
+ Expr *buildMemberExpr(Expr *Base, ValueDecl *Member) {
+ DeclAccessPair MemberDAP = DeclAccessPair::make(Member, AS_none);
+ MemberExpr *Result = SemaSYCLRef.SemaRef.BuildMemberExpr(
+ Base, /*IsArrow */ false, SrcLoc, NestedNameSpecifierLoc(), SrcLoc,
+ Member, MemberDAP,
+ /*HadMultipleCandidates*/ false,
+ DeclarationNameInfo(Member->getDeclName(), SrcLoc), Member->getType(),
+ VK_LValue, OK_Ordinary);
+ return Result;
+ }
+ Expr *buildDerivedToBaseCast(Expr *Base, CXXBaseSpecifier *BS) {
+ CXXCastPath BasePath;
+ QualType DerivedTy = SemaSYCLRef.getASTContext().getCanonicalTagType(CurRD);
+ QualType BaseTy = BS->getType();
+ SemaSYCLRef.SemaRef.CheckDerivedToBaseConversion(DerivedTy, BaseTy, SrcLoc,
+ SourceRange(), &BasePath,
+ /*IgnoreBaseAccess*/ true);
+ auto *Cast = ImplicitCastExpr::Create(
+ SemaSYCLRef.getASTContext(), BaseTy, CK_DerivedToBase, Base,
+ /* CXXCastPath=*/&BasePath, VK_LValue, FPOptionsOverride());
+ return Cast;
+ }
+
+ void createResultingArg() {
+ SmallVector<Expr *, 16> MemberExprBases;
+ auto *Param = cast<ValueDecl *>(ObjectAccessPath.front());
+ QualType ParamTy = Param->getType().getNonReferenceType();
+ Expr *Base =
+ SemaSYCLRef.SemaRef.BuildDeclRefExpr(Param, ParamTy, VK_LValue, SrcLoc);
+ MemberExprBases.push_back(Base);
+
+ for (auto Parent : ObjectAccessPath) {
+ if (auto *FD = Parent.dyn_cast<FieldDecl *>()) {
+ MemberExprBases.push_back(buildMemberExpr(MemberExprBases.back(), FD));
+ } else if (auto *BS = Parent.dyn_cast<CXXBaseSpecifier *>()) {
+ MemberExprBases.push_back(
+ buildDerivedToBaseCast(MemberExprBases.back(), BS));
+ }
+ }
+ ResultingArgs.push_back(MemberExprBases.back());
+ }
+
+public:
+ KernelSpecialParamsCreator(SemaSYCL &SR, SourceLocation Loc,
+ llvm::SmallVectorImpl<Expr *> &ResultingArgs)
+ : SubobjectVisitor<KernelSpecialParamsCreator>(SR.getASTContext()),
+ SemaSYCLRef(SR), SrcLoc(Loc), ResultingArgs(ResultingArgs) {}
+
+ void traverseRecord(RecordDecl *RD) {
+ CurRD = RD;
+ SubobjectVisitor::traverseRecord(RD);
+ }
+
+ void checkParameter(ValueDecl *PVD) {
+ ObjectAccessPath.push_back(PVD);
+ // Check the immediate type of the parameter.
+ visit(PVD->getType());
+ ObjectAccessPath.pop_back();
+ assert(ObjectAccessPath.empty());
+ }
+
+ bool visitBaseSpecifierPre(CXXBaseSpecifier *BS) {
+ ObjectAccessPath.push_back(BS);
+
+ // Do not visit inside of special types.
+ return !isSyclSpecialType(BS->getType());
+ }
+
+ bool visitFieldDeclPre(FieldDecl *FD) {
+ ObjectAccessPath.push_back(FD);
+
+ // Do not visit inside of special types.
+ return !isSyclSpecialType(FD->getType());
+ }
+ void visitFieldDeclPost(FieldDecl *FD) {
+ if (isSyclSpecialType(FD->getType()))
+ createResultingArg();
+
+ ObjectAccessPath.pop_back();
+ }
+
+ void visitBaseSpecifierPost(CXXBaseSpecifier *BS) {
+ // TODO: test bases
+ if (isSyclSpecialType(BS->getType()))
+ createResultingArg();
+
+ ObjectAccessPath.pop_back();
+ }
+};
+
+static void createArgumentsForSpecialTypes(SmallVectorImpl<Expr *> &Args,
+ ValueDecl *KernelArgObj,
+ SourceLocation Loc,
+ SemaSYCL &SemaSYCLRef) {
+ KernelSpecialParamsCreator KSPC(SemaSYCLRef, Loc, Args);
+ KSPC.checkParameter(KernelArgObj);
+}
+
// 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) {
+static 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
// that parameter references occur within the correct scope.
assert(SemaRef.CurContext == FD && "The current declaration context does not "
@@ -541,9 +657,9 @@ 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) {
+StmtResult BuildSYCLKernelLaunchCallStmt(
+ Sema &SemaRef, FunctionDecl *FD, const SYCLKernelInfo *SKI, Expr *IdExpr,
+ SourceLocation Loc, SmallVectorImpl<QualType> &SpecialArgTys) {
SmallVector<Stmt *> Stmts;
// IdExpr may be null if name lookup failed.
if (IdExpr) {
@@ -575,8 +691,39 @@ StmtResult BuildSYCLKernelLaunchCallStmt(Sema &SemaRef, FunctionDecl *FD,
SemaRef.BuildCallExpr(SemaRef.getCurScope(), IdExpr, Loc, Args, Loc);
if (LaunchResult.isInvalid())
return StmtError();
+ Expr *BaseForSubsCall =
+ SemaRef.MaybeCreateExprWithCleanups(LaunchResult).get();
+ if (!BaseForSubsCall->getType()->isVoidType()) {
+ // FIXME: diagnose that sycl_kernel_launch call returned a callable
+ // object. Default diagnostic here is very uncldear
+ llvm::SmallVector<Expr *, 12> SpecialArgs;
+ for (auto Param : FD->parameters()) {
+ if (Param->getType()->isRecordType())
+ createArgumentsForSpecialTypes(SpecialArgs, Param, Loc,
+ SemaRef.SYCL());
+ }
+ ExprResult Result = SemaRef.BuildCallExpr(
+ SemaRef.getCurScope(), BaseForSubsCall, Loc, SpecialArgs, Loc);
+ if (Result.isInvalid())
+ return StmtError();
+
+ // Now gather types for device code generation. Callable object returned
+ // by sycl_kernel_launch call returns type_list object whose template
+ // arguments describe types of additional kernel arguments required for
+ // special objects, i.e. SYCL accessors/samplers/streams etc.
+ QualType Ty = Result.get()->getType();
+ // FIXME: that also needs to be diagnosed somewhere.
+ auto *TST = Ty->getAs<TemplateSpecializationType>();
+ if (!TST)
+ return StmtError();
+ for (auto Arg : TST->template_arguments()) {
+ SpecialArgTys.push_back(Arg.getAsType().getCanonicalType());
+ }
- Stmts.push_back(SemaRef.MaybeCreateExprWithCleanups(LaunchResult).get());
+ Stmts.push_back(SemaRef.MaybeCreateExprWithCleanups(Result).get());
+ } else {
+ Stmts.push_back(BaseForSubsCall);
+ }
}
}
@@ -639,14 +786,21 @@ class OutlinedFunctionDeclBodyInstantiator
FunctionDecl *FD;
};
-OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
- FunctionDecl *FD,
- CompoundStmt *Body) {
+DeclResult
+BuildSYCLKernelEntryPointOutline(Sema &SemaRef, FunctionDecl *FD,
+ CompoundStmt *Body,
+ SmallVectorImpl<QualType> &SpecialArgTys,
+ Expr *IdExpr, SourceLocation Loc) {
using ParmDeclMap = OutlinedFunctionDeclBodyInstantiator::ParmDeclMap;
ParmDeclMap ParmMap;
OutlinedFunctionDecl *OFD = OutlinedFunctionDecl::Create(
- SemaRef.getASTContext(), FD, FD->getNumParams());
+ SemaRef.getASTContext(), FD, FD->getNumParams() + SpecialArgTys.size());
+
+ // CurContext is skep-attributed function but we're actually building device
+ // version of it which is a different DeclContext, so push it on the stack.
+ Sema::ContextRAII SavedContext(SemaRef, OFD);
+
unsigned i = 0;
for (ParmVarDecl *PVD : FD->parameters()) {
ImplicitParamDecl *IPD = ImplicitParamDecl::Create(
@@ -656,14 +810,88 @@ OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
ParmMap[PVD] = IPD;
++i;
}
-
OutlinedFunctionDeclBodyInstantiator OFDBodyInstantiator(SemaRef, ParmMap,
FD);
- Stmt *OFDBody = OFDBodyInstantiator.TransformStmt(Body).get();
+ Stmt *TransformedBody = OFDBodyInstantiator.TransformStmt(Body).get();
+
+ // Create kernel parameters for special types and create arguments to
+ // sycl_handle_special_kernel_parameters call.
+ // This is synthesizing the following pseudo-code:
+ // void kernel-entry-point(lambda-from-f kernelFunc, buffer_t* X, int Y) {
+ // sycl_handle_special_kernel_parameters<kernel-name-type>(kernelFunc.sout)(X, Y);
+ // {
+ // // This is copied body of the orignal skep-attributed function.
+ // kernelFunc();
+ // }
+ // }
+ // where sout is has type marked with sycl_special_kernel_parameter attribute.
+ Stmt *OFDBody;
+ if (IdExpr && !SpecialArgTys.empty()) {
+ SmallVector<Expr *, 8> HandleArgs;
+ for (unsigned I = 0; I < FD->getNumParams(); ++I) {
+ auto Param = OFD->getParam(I);
+ if (Param->getType()->isRecordType())
+ createArgumentsForSpecialTypes(HandleArgs, Param, Loc, SemaRef.SYCL());
+ }
+
+ // FIXME add better diagnosing.
+ // Sema::CodeSynthesisContext CSC;
+ // CSC.Kind =
+ // Sema::CodeSynthesisContext::SYCLKernelLaunchOverloadResolution;
+ // CSC.Entity = FD;
+ // CSC.CallArgs = Args.data();
+ // CSC.NumCallArgs = Args.size();
+ // Sema::ScopedCodeSynthesisContext ScopedCSC(SemaRef, CSC);
+
+ // Handle args for sycl_handle_special_kernel_parameters call, these are
+ // coming from subobjects with sycl_special_kernel_parameter attribute
+ // within skep-attributed function arguments, SpecialArgs are additional kernel
+ // arguments that are needed to initialize special subobjects and they go
+ // to the subsequent call.
+ SmallVector<Expr *, 12> SpecialArgs;
+ for (auto QT : SpecialArgTys) {
+ ImplicitParamDecl *IPD = ImplicitParamDecl::Create(
+ SemaRef.getASTContext(), OFD, SourceLocation(),
+ &SemaRef.getASTContext().Idents.get("idk"), QT,
+ ImplicitParamKind::Other);
+ OFD->setParam(i, IPD);
+ ++i;
+ ExprResult Arg =
+ SemaRef.BuildDeclRefExpr(IPD, QT, VK_LValue, SourceLocation());
+ assert(!Arg.isInvalid() && "synthesized code generation failed?");
+ SpecialArgs.push_back(Arg.get());
+ }
+
+ // This generates sycl_handle_special_kernel_parameters<kernel-name-type>(kernelFunc.sout)
+ ExprResult FirstHandleCallResult = SemaRef.BuildCallExpr(
+ SemaRef.getCurScope(), IdExpr, Loc, HandleArgs, Loc);
+ if (FirstHandleCallResult.isInvalid())
+ return true;
+ // FIXME: diagnose that sycl_special_kernel_parameter call returned a
+ // callable object. Default diagnostic here is very uncldear
+
+ Expr *BaseForSubsCall =
+ SemaRef.MaybeCreateExprWithCleanups(FirstHandleCallResult).get();
+ ExprResult Result = SemaRef.BuildCallExpr(
+ SemaRef.getCurScope(), BaseForSubsCall, Loc, SpecialArgs, Loc);
+ if (Result.isInvalid())
+ return true;
+
+ SmallVector<Stmt *> Stmts;
+ // Make sure to push kernel argument processing result first, before the
+ // transformed body of skep-attributed function.
+ Stmts.push_back(SemaRef.MaybeCreateExprWithCleanups(Result).get());
+ Stmts.push_back(TransformedBody);
+ OFDBody = CompoundStmt::Create(SemaRef.getASTContext(), Stmts,
+ FPOptionsOverride(), Loc, Loc);
+ } else {
+ OFDBody = TransformedBody;
+ }
+
OFD->setBody(OFDBody);
OFD->setNothrow();
-
return OFD;
+
}
class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
@@ -773,7 +1001,8 @@ bool verifyKernelParams(FunctionDecl *FD, SemaSYCL &SemaSYCLRef) {
StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl *FD,
CompoundStmt *Body,
- Expr *LaunchIdExpr) {
+ Expr *LaunchIdExpr,
+ Expr *HandleSpecParamsExpr) {
assert(!FD->isInvalidDecl());
assert(!FD->isTemplated());
assert(FD->hasPrototype());
@@ -797,27 +1026,31 @@ StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl *FD,
if (verifyKernelParams(FD, *this))
return StmtError();
- // Build the outline of the synthesized device entry point function.
- OutlinedFunctionDecl *OFD =
- BuildSYCLKernelEntryPointOutline(SemaRef, FD, Body);
- assert(OFD);
-
+ SourceLocation Loc = Body->getLBracLoc();
// Build the host kernel launch statement. An appropriate source location
// is required to emit diagnostics.
- SourceLocation Loc = Body->getLBracLoc();
- StmtResult LaunchResult =
- BuildSYCLKernelLaunchCallStmt(SemaRef, FD, &SKI, LaunchIdExpr, Loc);
+ llvm::SmallVector<QualType, 8> SpecialArgTys;
+ StmtResult LaunchResult = BuildSYCLKernelLaunchCallStmt(
+ SemaRef, FD, &SKI, LaunchIdExpr, Loc, SpecialArgTys);
+
if (LaunchResult.isInvalid())
return StmtError();
- Stmt *NewBody =
- new (getASTContext()) SYCLKernelCallStmt(Body, LaunchResult.get(), OFD);
+ // Build the outline of the synthesized device entry point function.
+ DeclResult OFD = BuildSYCLKernelEntryPointOutline(
+ SemaRef, FD, Body, SpecialArgTys, HandleSpecParamsExpr, Loc);
+
+ if (OFD.isInvalid())
+ return StmtError();
+
+ Stmt *NewBody = new (getASTContext()) SYCLKernelCallStmt(
+ Body, LaunchResult.get(), cast<OutlinedFunctionDecl>(OFD.get()));
return NewBody;
}
-StmtResult SemaSYCL::BuildUnresolvedSYCLKernelCallStmt(CompoundStmt *Body,
- Expr *LaunchIdExpr) {
- return UnresolvedSYCLKernelCallStmt::Create(SemaRef.getASTContext(), Body,
- LaunchIdExpr);
+StmtResult SemaSYCL::BuildUnresolvedSYCLKernelCallStmt(
+ CompoundStmt *Body, Expr *LaunchIdExpr, Expr *HandleSpecParamsExpr) {
+ return UnresolvedSYCLKernelCallStmt::Create(
+ SemaRef.getASTContext(), Body, LaunchIdExpr, HandleSpecParamsExpr);
}
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 3c8fcbe582b43..dc36a349146e5 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -13177,13 +13177,18 @@ StmtResult TreeTransform<Derived>::TransformUnresolvedSYCLKernelCallStmt(
if (IdExpr.isInvalid())
return StmtError();
+ ExprResult SpecArgsIdExpr =
+ getDerived().TransformExpr(S->getSpecArgsIdExpr());
+ if (SpecArgsIdExpr.isInvalid())
+ return StmtError();
+
StmtResult Body = getDerived().TransformStmt(S->getOriginalStmt());
if (Body.isInvalid())
return StmtError();
StmtResult SR = SemaRef.SYCL().BuildSYCLKernelCallStmt(
cast<FunctionDecl>(SemaRef.CurContext), cast<CompoundStmt>(Body.get()),
- IdExpr.get());
+ IdExpr.get(), SpecArgsIdExpr.get());
if (SR.isInvalid())
return StmtError();
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..8857518856c57 100644
--- a/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp
+++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp
@@ -37,6 +37,9 @@ template<int> struct K {
template<typename KernelName, typename... Ts>
void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KernelName, typename... Ts>
+void sycl_handle_special_kernel_parameters(Ts...) {}
+
[[clang::sycl_kernel_entry_point(KN<1>)]]
void skep1() {
}
diff --git a/clang/test/ASTSYCL/ast-dump-sycl-kernel-decomposition.cpp b/clang/test/ASTSYCL/ast-dump-sycl-kernel-decomposition.cpp
new file mode 100644
index 0000000000000..67e7b0266dc32
--- /dev/null
+++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-decomposition.cpp
@@ -0,0 +1,141 @@
+// Tests without serialization:
+// RUN: %clang_cc1 -std=c++17 -triple spirv64-unknown-unknown -fsycl-is-device \
+// RUN: -ast-dump %s \
+// RUN: | FileCheck %s
+// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-unknown -fsycl-is-host \
+// RUN: -ast-dump %s \
+// RUN: | FileCheck %s
+
+// Thes test validates the AST body produced for functions declared with the
+// sycl_kernel_entry_point attribute in case an argument of such function
+// contains an object that requires decomposition.
+
+// CHECK: TranslationUnitDecl {{.*}}
+
+// A unique kernel name type is required for each declared kernel entry point.
+template<int> struct KN;
+
+struct [[clang::sycl_special_kernel_parameter]] EmptySpecial {
+ int data;
+};
+
+template<typename T>
+struct Wrapper {
+ T data;
+ int *data1;
+};
+
+template<typename T>
+auto set_kernel_arg(const T &t) {
+ return t;
+}
+
+auto set_kernel_arg(EmptySpecial &a) {
+ return a.data;
+}
+
+template<typename KernelName, typename... Ts>
+auto sycl_handle_special_kernel_parameters(Ts...) {
+ return [](auto ...Args){ return; };
+}
+
+template<typename... Ts>
+struct type_list {};
+
+template <typename KernelName, typename... Ts>
+auto sycl_kernel_launch(const char *, Ts...) {
+
+ return [&](auto&&... extra_host_args) {
+ return type_list<decltype(set_kernel_arg(extra_host_args))...>{};
+ };
+}
+
+
+template <typename KN, typename KT>
+[[clang::sycl_kernel_entry_point(KN)]] void k(KT Kernel) {
+ Kernel();
+}
+// CHECK: |-FunctionTemplateDecl {{.*}} k{{.*}}
+// CHECK-NEXT: | |-TemplateTypeParmDecl {{.*}} referenced typename depth 0 index 0 KN
+// CHECK-NEXT: | |-TemplateTypeParmDecl {{.*}} referenced typename depth 0 index 1 KT
+// CHECK-NEXT: | |-FunctionDecl {{.*}} k 'void (KT)'
+// CHECK-NEXT: | | |-ParmVarDecl {{.*}} referenced Kernel 'KT'
+// CHECK-NEXT: | | |-UnresolvedSYCLKernelCallStmt {{.*}}
+// CHECK-NEXT: | | | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | | `-CallExpr {{.*}} '<dependent type>'
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'Kernel' 'KT'
+// CHECK-NEXT: | | | |-UnresolvedLookupExpr {{.*}} '<dependent type>' lvalue (ADL) = 'sycl_kernel_launch' {{.*}}
+// CHECK-NEXT: | | | | `-TemplateArgument type 'KN':'type-parameter-0-0'
+// CHECK-NEXT: | | | | `-TemplateTypeParmType {{.*}} 'KN' dependent depth 0 index 0
+// CHECK-NEXT: | | | | `-TemplateTypeParm {{.*}} 'KN'
+// CHECK-NEXT: | | | `-UnresolvedLookupExpr {{.*}} '<dependent type>' lvalue (ADL) = 'sycl_handle_special_kernel_parameters' {{.*}}
+// CHECK-NEXT: | | | `-TemplateArgument type 'KN':'type-parameter-0-0'
+// CHECK-NEXT: | | | `-TemplateTypeParmType {{.*}} 'KN' dependent depth 0 index 0
+// CHECK-NEXT: | | | `-TemplateTypeParm {{.*}} 'KN'
+// CHECK-NEXT: | | `-SYCLKernelEntryPointAttr {{.*}} KN
+// CHECK-NEXT: | `-FunctionDecl {{.*}} used k {{.*}} implicit_instantiation instantiated_from {{.*}}
+// CHECK-NEXT: | |-TemplateArgument type 'KN<0>'
+// CHECK-NEXT: | | `-RecordType {{.*}} 'KN<0>' canonical
+// CHECK-NEXT: | | `-ClassTemplateSpecialization {{.*}} 'KN'
+// CHECK-NEXT: | |-TemplateArgument type '{{.*}}'
+// CHECK-NEXT: | | `-RecordType {{.*}} canonical
+// CHECK-NEXT: | | `-CXXRecord {{.*}}
+// CHECK-NEXT: | |-ParmVarDecl {{.*}} used Kernel {{.*}}
+// CHECK-NEXT: | |-SYCLKernelCallStmt {{.*}}
+// CHECK-NEXT: | | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | `-CXXOperatorCallExpr {{.*}} 'void' '()'
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)() const' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod {{.*}} 'operator()' 'void () const'
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const {{.*}}' lvalue <NoOp>
+// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} lvalue ParmVar {{.*}} 'Kernel' {{.*}}
+// CHECK-NEXT: | | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | `-ExprWithCleanups {{.*}} 'type_list<{{.*}}>'
+// CHECK-NEXT: | | | `-CXXOperatorCallExpr {{.*}} 'type_list<{{.*}}>' '()'
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'type_list<{{.*}}> (*)(EmptySpecial &) const' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'type_list<{{.*}}> (EmptySpecial &) const' lvalue CXXMethod {{.*}} 'operator()' '{{.*}}'
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const {{.*}}' lvalue <NoOp>
+// CHECK-NEXT: | | | | `-MaterializeTemporaryExpr {{.*}} '{{.*}}' lvalue
+// CHECK-NEXT: | | | | `-CallExpr {{.*}} '{{.*}}'
+// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} '{{.*}}' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} '{{.*}}' lvalue Function {{.*}} 'sycl_kernel_launch' {{.*}}
+// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: | | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi0EE"
+// CHECK-NEXT: | | | | `-CXXConstructExpr {{.*}} '{{.*}}' 'void ({{.*}} &&) noexcept'
+// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} '{{.*}}' xvalue <NoOp>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} lvalue ParmVar {{.*}} 'Kernel' {{.*}}
+// CHECK-NEXT: | | | `-MemberExpr {{.*}} 'EmptySpecial' lvalue .data {{.*}}
+// CHECK-NEXT: | | | `-MemberExpr {{.*}} 'Wrapper<EmptySpecial>' lvalue . {{.*}}
+// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} lvalue ParmVar {{.*}} 'Kernel' {{.*}}
+// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
+// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used Kernel {{.*}}
+// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used idk {{.*}}
+// CHECK-NEXT: | | `-CompoundStmt {{.*}}
+// CHECK-NEXT: | | |-ExprWithCleanups {{.*}} 'void'
+// CHECK-NEXT: | | | `-CXXOperatorCallExpr {{.*}} 'void' '()'
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(int) const' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (int) const' lvalue CXXMethod {{.*}} 'operator()' '{{.*}}'
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const {{.*}}' lvalue <NoOp>
+// CHECK-NEXT: | | | | `-MaterializeTemporaryExpr {{.*}} '{{.*}}' lvalue
+// CHECK-NEXT: | | | | `-CallExpr {{.*}} '{{.*}}'
+// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} '{{.*}}' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} '{{.*}}' lvalue Function {{.*}} 'sycl_handle_special_kernel_parameters' {{.*}}
+// CHECK-NEXT: | | | | `-CXXConstructExpr {{.*}} 'EmptySpecial' 'void (const EmptySpecial &) noexcept'
+// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'const EmptySpecial' lvalue <NoOp>
+// CHECK-NEXT: | | | | `-MemberExpr {{.*}} 'EmptySpecial' lvalue .data {{.*}}
+// CHECK-NEXT: | | | | `-MemberExpr {{.*}} 'Wrapper<EmptySpecial>' lvalue . {{.*}}
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} lvalue ImplicitParam {{.*}} 'Kernel' {{.*}}
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'int' lvalue ImplicitParam {{.*}} 'idk' 'int'
+// CHECK-NEXT: | | `-CompoundStmt {{.*}}
+// CHECK-NEXT: | | `-CXXOperatorCallExpr {{.*}} 'void' '()'
+// CHECK-NEXT: | | |-ImplicitCastExpr {{.*}} 'void (*)() const' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod {{.*}} 'operator()' 'void () const'
+// CHECK-NEXT: | | `-ImplicitCastExpr {{.*}} 'const {{.*}}' lvalue <NoOp>
+// CHECK-NEXT: | | `-DeclRefExpr {{.*}} lvalue ImplicitParam {{.*}} 'Kernel' {{.*}}
+// CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} struct KN<0>
+
+void case1() {
+ Wrapper<EmptySpecial> KernelArg;
+ k<KN<0>>([KernelArg](){});
+}
+// CHECK: `-FunctionDecl {{.*}} case1 'void ()'
diff --git a/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp b/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
index cca10365e7686..5f88d87894a08 100644
--- a/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
+++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
@@ -31,6 +31,9 @@ template<int, int=0> struct KN;
template<typename KernelName, typename... Ts>
void sycl_kernel_launch(const char *, Ts... Args) {}
+template<typename KernelName, typename... Ts>
+void sycl_handle_special_kernel_parameters(Ts...) {}
+
[[clang::sycl_kernel_entry_point(KN<1>)]]
void skep1() {
}
diff --git a/clang/test/ASTSYCL/ast-print-sycl-kernel-call.cpp b/clang/test/ASTSYCL/ast-print-sycl-kernel-call.cpp
index 5adaa367ed9c1..f32b931205042 100644
--- a/clang/test/ASTSYCL/ast-print-sycl-kernel-call.cpp
+++ b/clang/test/ASTSYCL/ast-print-sycl-kernel-call.cpp
@@ -1,6 +1,9 @@
// RUN: %clang_cc1 -fsycl-is-host -ast-print %s -o - | FileCheck %s
// RUN: %clang_cc1 -fsycl-is-device -ast-print %s -o - | FileCheck %s
+template<typename KernelName, typename... Ts>
+void sycl_handle_special_kernel_parameters(Ts...) {}
+
struct sycl_kernel_launcher {
template<typename KernelName, typename... Ts>
void sycl_kernel_launch(const char *, Ts...) {}
diff --git a/clang/test/CodeGenSYCL/function-attrs.cpp b/clang/test/CodeGenSYCL/function-attrs.cpp
index 60d3cf10055ec..efb0081484522 100644
--- a/clang/test/CodeGenSYCL/function-attrs.cpp
+++ b/clang/test/CodeGenSYCL/function-attrs.cpp
@@ -29,6 +29,9 @@ int foo() {
template <typename Name, typename... Ts>
void sycl_kernel_launch(Ts...) {}
+template <typename Name, typename... Ts>
+void sycl_handle_special_kernel_parameters(Ts...) {}
+
template <typename Name, typename Func>
[[clang::sycl_kernel_entry_point(Name)]] void kernel_single_task(const Func &kernelFunc) {
kernelFunc();
diff --git a/clang/test/CodeGenSYCL/kernel-arg-decomposition.cpp b/clang/test/CodeGenSYCL/kernel-arg-decomposition.cpp
new file mode 100644
index 0000000000000..b7aa38cc9a1cb
--- /dev/null
+++ b/clang/test/CodeGenSYCL/kernel-arg-decomposition.cpp
@@ -0,0 +1,96 @@
+// RUN: %clang_cc1 -fsycl-is-host -emit-llvm -triple x86_64-unknown-linux-gnu -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-HOST %s
+// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -aux-triple x86_64-unknown-linux-gnu -triple spirv64-unknown-unknown -std=c++17 %s -o - | FileCheck --check-prefixes=CHECK-DEVICE %s
+
+// A unique kernel name type is required for each declared kernel entry point.
+template<int> struct KN;
+
+struct [[clang::sycl_special_kernel_parameter]] EmptySpecial {
+ int data;
+};
+
+template<typename T>
+struct Wrapper {
+ T data;
+ int *data1;
+};
+
+template<typename T>
+auto set_kernel_arg(const T &t) {
+ return t;
+}
+
+auto set_kernel_arg(EmptySpecial &a) {
+ return a.data;
+}
+
+template<typename KernelName, typename... Ts>
+auto sycl_handle_special_kernel_parameters(Ts...) {
+ return [](auto ...Args){ return; };
+}
+
+template<typename... Ts>
+struct type_list {};
+
+template <typename KernelName, typename... Ts>
+auto sycl_kernel_launch(const char *, Ts...) {
+
+ return [&](auto&&... extra_host_args) {
+ return type_list<decltype(set_kernel_arg(extra_host_args))...>{};
+ };
+}
+
+
+template <typename KN, typename KT>
+[[clang::sycl_kernel_entry_point(KN)]] void kernel_entry_point(KT Kernel) {
+ Kernel();
+}
+
+void case1() {
+ Wrapper<EmptySpecial> KernelArg;
+ kernel_entry_point<KN<0>>([KernelArg](){});
+}
+
+// CHECK-HOST-LABEL: define internal void @_Z18kernel_entry_pointI2KNILi0EEZ5case1vEUlvE_EvT0_(
+// CHECK-HOST-SAME: i32 [[KERNEL_COERCE0:%.*]], ptr [[KERNEL_COERCE1:%.*]])
+// CHECK-HOST: [[ENTRY:.*:]]
+// CHECK-HOST-NEXT: [[KERNEL:%.*]] = alloca [[CLASS_ANON:%.*]], align 8
+// CHECK-HOST-NEXT: [[REF_TMP:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 1
+// CHECK-HOST-NEXT: [[AGG_TMP:%.*]] = alloca [[CLASS_ANON]], align 8
+// CHECK-HOST-NEXT: [[UNDEF_AGG_TMP:%.*]] = alloca [[CLASS_ANON_0]], align 1
+// CHECK-HOST-NEXT: [[UNDEF_AGG_TMP1:%.*]] = alloca [[STRUCT_TYPE_LIST:%.*]], align 1
+// CHECK-HOST-NEXT: [[TMP0:%.*]] = getelementptr inbounds nuw { i32, ptr }, ptr [[KERNEL]], i32 0, i32 0
+// CHECK-HOST-NEXT: store i32 [[KERNEL_COERCE0]], ptr [[TMP0]], align 8
+// CHECK-HOST-NEXT: [[TMP1:%.*]] = getelementptr inbounds nuw { i32, ptr }, ptr [[KERNEL]], i32 0, i32 1
+// CHECK-HOST-NEXT: store ptr [[KERNEL_COERCE1]], ptr [[TMP1]], align 8
+// CHECK-HOST-NEXT: call void @llvm.memcpy.p0.p0.i64(ptr align 8 [[AGG_TMP]], ptr align 8 [[KERNEL]], i64 16, i1 false)
+// CHECK-HOST-NEXT: [[TMP2:%.*]] = getelementptr inbounds nuw { i32, ptr }, ptr [[AGG_TMP]], i32 0, i32 0
+// CHECK-HOST-NEXT: [[TMP3:%.*]] = load i32, ptr [[TMP2]], align 8
+// CHECK-HOST-NEXT: [[TMP4:%.*]] = getelementptr inbounds nuw { i32, ptr }, ptr [[AGG_TMP]], i32 0, i32 1
+// CHECK-HOST-NEXT: [[TMP5:%.*]] = load ptr, ptr [[TMP4]], align 8
+// CHECK-HOST-NEXT: call void @_Z18sycl_kernel_launchI2KNILi0EEJZ5case1vEUlvE_EEDaPKcDpT0_(ptr noundef @.str, i32 [[TMP3]], ptr [[TMP5]])
+// CHECK-HOST-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[CLASS_ANON]], ptr [[KERNEL]], i32 0, i32 0
+// CHECK-HOST-NEXT: [[DATA:%.*]] = getelementptr inbounds nuw [[STRUCT_WRAPPER:%.*]], ptr [[TMP6]], i32 0, i32 0
+// CHECK-HOST-NEXT: call void @_ZZ18sycl_kernel_launchI2KNILi0EEJZ5case1vEUlvE_EEDaPKcDpT0_ENKUlDpOT_E_clIJR12EmptySpecialEEEDaS9_(ptr noundef nonnull align 1 dereferenceable(1) [[REF_TMP]], ptr noundef nonnull align 4 dereferenceable(4) [[DATA]])
+// CHECK-HOST-NEXT: ret void
+
+// CHECK-DEVICE: define spir_kernel void @_ZTS2KNILi0EE(ptr noundef byval(%class.anon) align 8 [[KERNEL:%.*]], i32 noundef [[IDK:%.*]])
+// CHECK-DEVICE: [[ENTRY:.*:]]
+// CHECK-DEVICE: [[IDK_ADDR:%.*]] = alloca i32, align 4
+// CHECK-DEVICE: [[REF_TMP:%.*]] = alloca [[CLASS_ANON_0:%.*]], align 1
+// CHECK-DEVICE: [[AGG_TMP:%.*]] = alloca [[STRUCT_EMPTYSPECIAL:%.*]], align 4
+// CHECK-DEVICE: [[IDK_ADDR_ASCAST:%.*]] = addrspacecast ptr [[IDK_ADDR]] to ptr addrspace(4)
+// CHECK-DEVICE: [[REF_TMP_ASCAST:%.*]] = addrspacecast ptr [[REF_TMP]] to ptr addrspace(4)
+// CHECK-DEVICE: [[AGG_TMP_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP]] to ptr addrspace(4)
+// CHECK-DEVICE: [[KERNEL_ASCAST:%.*]] = addrspacecast ptr [[KERNEL]] to ptr addrspace(4)
+// CHECK-DEVICE: store i32 [[IDK]], ptr addrspace(4) [[IDK_ADDR_ASCAST]], align 4
+// CHECK-DEVICE: [[REF_TMP_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[REF_TMP_ASCAST]] to ptr
+// CHECK-DEVICE: [[TMP0:%.*]] = getelementptr inbounds nuw [[CLASS_ANON:%.*]], ptr addrspace(4) [[KERNEL_ASCAST]], i32 0, i32 0
+// CHECK-DEVICE: [[DATA:%.*]] = getelementptr inbounds nuw [[STRUCT_WRAPPER:%.*]], ptr addrspace(4) [[TMP0]], i32 0, i32 0
+// CHECK-DEVICE: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP_ASCAST]], ptr addrspace(4) align 8 [[DATA]], i64 4, i1 false)
+// CHECK-DEVICE: [[AGG_TMP_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP_ASCAST]] to ptr
+// CHECK-DEVICE: call spir_func void @_Z37sycl_handle_special_kernel_parametersI2KNILi0EEJ12EmptySpecialEEDaDpT0_(ptr dead_on_unwind writable sret([[CLASS_ANON_0]]) align 1 [[REF_TMP_ASCAST_ASCAST]], ptr noundef byval([[STRUCT_EMPTYSPECIAL]]) align 4 [[AGG_TMP_ASCAST_ASCAST]])
+// CHECK-DEVICE: [[TMP1:%.*]] = load i32, ptr addrspace(4) [[IDK_ADDR_ASCAST]], align 4
+// CHECK-DEVICE: call spir_func void @_ZZ37sycl_handle_special_kernel_parametersI2KNILi0EEJ12EmptySpecialEEDaDpT0_ENKUlDpT_E_clIJiEEEDaS6_(ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) [[REF_TMP_ASCAST]], i32 noundef [[TMP1]])
+// CHECK-DEVICE: call spir_func void @_ZZ5case1vENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) [[KERNEL_ASCAST]])
+// CHECK-DEVICE: ret void
+
diff --git a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
index 270671284c3d1..a60ddcbeaa94e 100644
--- a/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
+++ b/clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp
@@ -33,6 +33,9 @@
template <typename KernelName, typename... Ts>
void sycl_kernel_launch(const char *, Ts...) {}
+template <typename KernelName, typename... Ts>
+void sycl_handle_special_kernel_parameters(Ts...) {}
+
struct single_purpose_kernel_name;
struct single_purpose_kernel {
void operator()() const {}
diff --git a/clang/test/CodeGenSYCL/sycl-kernel-entry-point-exceptions.cpp b/clang/test/CodeGenSYCL/sycl-kernel-entry-point-exceptions.cpp
index 8fe7a148a2f61..56137b934eeca 100644
--- a/clang/test/CodeGenSYCL/sycl-kernel-entry-point-exceptions.cpp
+++ b/clang/test/CodeGenSYCL/sycl-kernel-entry-point-exceptions.cpp
@@ -16,6 +16,9 @@ struct KT {
};
+template<typename KN, typename... Ts>
+void sycl_handle_special_kernel_parameters(Ts...) {}
+
// Validate that exception handling instructions are omitted when a
// potentially throwing sycl_kernel_entry_point attributed function
// calls a potentially throwing sycl_kernel_launch function (a thrown
diff --git a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
index c298593e2f1ab..be4fa5482552a 100644
--- a/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
+++ b/clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp
@@ -4,6 +4,9 @@
template<typename KN, typename... Ts>
void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KN, typename... Ts>
+void sycl_handle_special_kernel_parameters(Ts...) {}
+
template<typename KN, typename Func>
[[clang::sycl_kernel_entry_point(KN)]] void kernel(Func F){
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..33d2ac9a13aec 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp
@@ -44,6 +44,9 @@ template<int, int = 0> struct KN;
template<typename KNT, typename... Ts>
void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KNT, typename... Ts>
+void sycl_handle_special_kernel_parameters(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..0907850ca8d3f 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
@@ -20,6 +20,9 @@ struct type_info {
template<typename KernelName, typename... Ts>
void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KernelName, typename... Ts>
+void sycl_handle_special_kernel_parameters(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..346938039de1c 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
@@ -14,6 +14,9 @@ template<int N> using TTA = ST<N>; // #TTA-decl
template<typename KN, typename... Ts>
void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KN, typename... Ts>
+void sycl_handle_special_kernel_parameters(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..a0cbbfad420d7 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
@@ -21,6 +21,9 @@ template<int> struct KN;
template<typename KN, typename... Ts>
void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KN, typename... Ts>
+void sycl_handle_special_kernel_parameters(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..047c6268cfe6c 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
@@ -19,6 +19,9 @@ template<int> struct KN;
template<typename KN, typename... Ts>
void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KN, typename... Ts>
+void sycl_handle_special_kernel_parameters(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..e2a20f7960377 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
@@ -14,6 +14,9 @@ struct S1;
template<typename KernelName, typename... Ts>
void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KernelName, typename... Ts>
+void sycl_handle_special_kernel_parameters(Ts...) {}
+
// expected-warning at +3 {{redundant 'clang::sycl_kernel_entry_point' attribute}}
// expected-note at +1 {{previous attribute is here}}
[[clang::sycl_kernel_entry_point(S1),
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp
index 654564f675e54..5ac71f3dbc306 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp
@@ -14,6 +14,9 @@
template<typename KernelName, typename... Ts>
void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KernelName, typename... Ts>
+void sycl_handle_special_kernel_parameters(Ts...) {}
+
// FIXME: C++23 [temp.expl.spec]p12 states:
// FIXME: ... Similarly, attributes appearing in the declaration of a template
// FIXME: have no effect on an explicit specialization of that template.
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp
index ac2d5c1541528..b34ad0185c71c 100644
--- a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp
@@ -26,6 +26,9 @@ struct type_info {
template<typename KernelName, typename... Ts>
void sycl_kernel_launch(const char *, Ts...) {}
+template<typename KernelName, typename... Ts>
+void sycl_handle_special_kernel_parameters(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..59930e1b502b1 100644
--- a/clang/test/SemaSYCL/sycl-kernel-launch-ms-compat.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-launch-ms-compat.cpp
@@ -19,6 +19,9 @@ struct KT {
};
+template<typename KN, typename... Ts>
+void sycl_handle_special_kernel_parameters(Ts...) {}
+
namespace ok1 {
template<typename Derived>
struct base_handler {
diff --git a/clang/test/SemaSYCL/sycl-kernel-launch.cpp b/clang/test/SemaSYCL/sycl-kernel-launch.cpp
index 06452de35b3d2..9734fd9cd99b2 100644
--- a/clang/test/SemaSYCL/sycl-kernel-launch.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-launch.cpp
@@ -24,6 +24,9 @@ struct KT {
};
+template<typename KN, typename... Ts>
+void sycl_handle_special_kernel_parameters(Ts...) {}
+
// sycl_kernel_launch as function template at namespace scope.
namespace ok1 {
template<typename KN, typename... Ts>
>From 77b8bd3e14ee32aeec3c394c7c543830619b2991 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Wed, 8 Jul 2026 06:17:26 -0700
Subject: [PATCH 02/11] Fix serialization of UnresolvedSYCLKernelCallStmt
---
clang/lib/Serialization/ASTReaderStmt.cpp | 1 +
clang/lib/Serialization/ASTWriterStmt.cpp | 1 +
2 files changed, 2 insertions(+)
diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp
index 02ccf8d4d41c2..0563f51af6d87 100644
--- a/clang/lib/Serialization/ASTReaderStmt.cpp
+++ b/clang/lib/Serialization/ASTReaderStmt.cpp
@@ -619,6 +619,7 @@ void ASTStmtReader::VisitUnresolvedSYCLKernelCallStmt(
S->setOriginalStmt(cast<CompoundStmt>(Record.readSubStmt()));
S->setKernelLaunchIdExpr(Record.readExpr());
+ S->setSpecArgsIdExpr(Record.readExpr());
}
void ASTStmtReader::VisitPredefinedExpr(PredefinedExpr *E) {
diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp
index 4e9eadd730a56..dc5f56e36ae32 100644
--- a/clang/lib/Serialization/ASTWriterStmt.cpp
+++ b/clang/lib/Serialization/ASTWriterStmt.cpp
@@ -703,6 +703,7 @@ void ASTStmtWriter::VisitUnresolvedSYCLKernelCallStmt(
Record.AddStmt(S->getOriginalStmt());
Record.AddStmt(S->getKernelLaunchIdExpr());
+ Record.AddStmt(S->getSpecArgsIdExpr());
Code = serialization::STMT_UNRESOLVED_SYCL_KERNEL_CALL;
}
>From ec657652b56bc70d4c83c2d6499685ae4d9c1c86 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Tue, 14 Jul 2026 03:44:21 -0700
Subject: [PATCH 03/11] Fix existing ast-dump-sycl-kernel-call-stmt.cpp test
---
clang/include/clang/AST/StmtSYCL.h | 3 +
clang/lib/Sema/SemaSYCL.cpp | 91 +++++++++++--------
.../ast-dump-sycl-kernel-call-stmt.cpp | 18 +++-
3 files changed, 70 insertions(+), 42 deletions(-)
diff --git a/clang/include/clang/AST/StmtSYCL.h b/clang/include/clang/AST/StmtSYCL.h
index cd682f4cea594..2c504e69c04f7 100644
--- a/clang/include/clang/AST/StmtSYCL.h
+++ b/clang/include/clang/AST/StmtSYCL.h
@@ -122,6 +122,9 @@ class UnresolvedSYCLKernelCallStmt : public Stmt {
void setOriginalStmt(CompoundStmt *CS) { OriginalStmt = CS; }
void setKernelLaunchIdExpr(Expr *IdExpr) { KernelLaunchIdExpr = IdExpr; }
+ void setSpecArgsIdExpr(Expr *IdExpr) {
+ HandleSYCLSpecialParamsIdExpr = IdExpr;
+ }
public:
static UnresolvedSYCLKernelCallStmt *Create(const ASTContext &C,
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 62d9aa653300b..ec62104553406 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -739,7 +739,7 @@ StmtResult BuildSYCLKernelLaunchCallStmt(
class OutlinedFunctionDeclBodyInstantiator
: public TreeTransform<OutlinedFunctionDeclBodyInstantiator> {
public:
- using ParmDeclMap = llvm::DenseMap<ParmVarDecl *, VarDecl *>;
+ using ParmDeclMap = llvm::DenseMap<VarDecl *, VarDecl *>;
OutlinedFunctionDeclBodyInstantiator(Sema &S, ParmDeclMap &M,
FunctionDecl *FD)
@@ -749,17 +749,17 @@ class OutlinedFunctionDeclBodyInstantiator
// A new set of AST nodes is always required.
bool AlwaysRebuild() { return true; }
- // Transform ParmVarDecl references to the supplied replacement variables.
+ // Transform VarDecl references to the supplied replacement variables.
ExprResult TransformDeclRefExpr(DeclRefExpr *DRE) {
- const ParmVarDecl *PVD = dyn_cast<ParmVarDecl>(DRE->getDecl());
- if (PVD) {
- ParmDeclMap::iterator I = MapRef.find(PVD);
+ const VarDecl *OrigVD = dyn_cast<VarDecl>(DRE->getDecl());
+ if (OrigVD) {
+ ParmDeclMap::iterator I = MapRef.find(OrigVD);
if (I != MapRef.end()) {
VarDecl *VD = I->second;
assert(SemaRef.getASTContext().hasSameUnqualifiedType(
- PVD->getType().getNonReferenceType(), VD->getType()));
+ OrigVD->getType().getNonReferenceType(), VD->getType()));
assert(!VD->getType().isMoreQualifiedThan(
- PVD->getType().getNonReferenceType(), SemaRef.getASTContext()));
+ OrigVD->getType().getNonReferenceType(), SemaRef.getASTContext()));
VD->setIsUsed();
return DeclRefExpr::Create(
SemaRef.getASTContext(), DRE->getQualifierLoc(),
@@ -797,10 +797,9 @@ BuildSYCLKernelEntryPointOutline(Sema &SemaRef, FunctionDecl *FD,
OutlinedFunctionDecl *OFD = OutlinedFunctionDecl::Create(
SemaRef.getASTContext(), FD, FD->getNumParams() + SpecialArgTys.size());
- // CurContext is skep-attributed function but we're actually building device
- // version of it which is a different DeclContext, so push it on the stack.
- Sema::ContextRAII SavedContext(SemaRef, OFD);
-
+ // We create everything in DeclContext of the original function (FD) and then
+ // replace all references to original function's parameters or variables
+ // using a TreeTransform.
unsigned i = 0;
for (ParmVarDecl *PVD : FD->parameters()) {
ImplicitParamDecl *IPD = ImplicitParamDecl::Create(
@@ -810,28 +809,28 @@ BuildSYCLKernelEntryPointOutline(Sema &SemaRef, FunctionDecl *FD,
ParmMap[PVD] = IPD;
++i;
}
- OutlinedFunctionDeclBodyInstantiator OFDBodyInstantiator(SemaRef, ParmMap,
- FD);
- Stmt *TransformedBody = OFDBodyInstantiator.TransformStmt(Body).get();
-
// Create kernel parameters for special types and create arguments to
// sycl_handle_special_kernel_parameters call.
// This is synthesizing the following pseudo-code:
// void kernel-entry-point(lambda-from-f kernelFunc, buffer_t* X, int Y) {
- // sycl_handle_special_kernel_parameters<kernel-name-type>(kernelFunc.sout)(X, Y);
+ // sycl_handle_special_kernel_parameters<kernel-name-type>(kernelFunc.sout).
+ // (X, Y);
// {
// // This is copied body of the orignal skep-attributed function.
// kernelFunc();
// }
// }
// where sout is has type marked with sycl_special_kernel_parameter attribute.
- Stmt *OFDBody;
+ Stmt *ModifiedBody;
if (IdExpr && !SpecialArgTys.empty()) {
+ // HandleArgs contains arguments for sycl_handle_special_kernel_parameters
+ // call, these are coming from subobjects of object whose type is marked
+ // with sycl_special_kernel_parameter attribute within skep-attributed
+ // function, i.e. kernelFunc.sout in the pseudo code above.
SmallVector<Expr *, 8> HandleArgs;
- for (unsigned I = 0; I < FD->getNumParams(); ++I) {
- auto Param = OFD->getParam(I);
- if (Param->getType()->isRecordType())
- createArgumentsForSpecialTypes(HandleArgs, Param, Loc, SemaRef.SYCL());
+ for (ParmVarDecl *PVD : FD->parameters()) {
+ if (PVD->getType()->isRecordType())
+ createArgumentsForSpecialTypes(HandleArgs, PVD, Loc, SemaRef.SYCL());
}
// FIXME add better diagnosing.
@@ -843,26 +842,34 @@ BuildSYCLKernelEntryPointOutline(Sema &SemaRef, FunctionDecl *FD,
// CSC.NumCallArgs = Args.size();
// Sema::ScopedCodeSynthesisContext ScopedCSC(SemaRef, CSC);
- // Handle args for sycl_handle_special_kernel_parameters call, these are
- // coming from subobjects with sycl_special_kernel_parameter attribute
- // within skep-attributed function arguments, SpecialArgs are additional kernel
- // arguments that are needed to initialize special subobjects and they go
- // to the subsequent call.
+ // SpecialArgs are additional kernel arguments that are needed to
+ // initialize special subobjects and they go to the subsequent call.
+ // These are (X, Y) in the pseudo code above.
SmallVector<Expr *, 12> SpecialArgs;
for (auto QT : SpecialArgTys) {
+ // Since these parameters do not exist in skep attributed function, we
+ // declare local variables instead.
+ VarDecl *VD = VarDecl::Create(
+ SemaRef.getASTContext(), SemaRef.getCurContext(), Loc, Loc,
+ &SemaRef.getASTContext().Idents.get("idk"), QT,
+ SemaRef.getASTContext().getTrivialTypeSourceInfo(QT, Loc), SC_None);
+ ExprResult Arg =
+ SemaRef.BuildDeclRefExpr(VD, QT, VK_LValue, SourceLocation());
+ assert(!Arg.isInvalid() && "synthesized code generation failed?");
+ SpecialArgs.push_back(Arg.get());
+
+ // Also create an implicit parameter for the device function.
ImplicitParamDecl *IPD = ImplicitParamDecl::Create(
SemaRef.getASTContext(), OFD, SourceLocation(),
&SemaRef.getASTContext().Idents.get("idk"), QT,
ImplicitParamKind::Other);
OFD->setParam(i, IPD);
++i;
- ExprResult Arg =
- SemaRef.BuildDeclRefExpr(IPD, QT, VK_LValue, SourceLocation());
- assert(!Arg.isInvalid() && "synthesized code generation failed?");
- SpecialArgs.push_back(Arg.get());
+ ParmMap[VD] = IPD;
}
- // This generates sycl_handle_special_kernel_parameters<kernel-name-type>(kernelFunc.sout)
+ // This generates
+ // sycl_handle_special_kernel_parameters<kernel-name-type>(kernelFunc.sout)
ExprResult FirstHandleCallResult = SemaRef.BuildCallExpr(
SemaRef.getCurScope(), IdExpr, Loc, HandleArgs, Loc);
if (FirstHandleCallResult.isInvalid())
@@ -872,26 +879,32 @@ BuildSYCLKernelEntryPointOutline(Sema &SemaRef, FunctionDecl *FD,
Expr *BaseForSubsCall =
SemaRef.MaybeCreateExprWithCleanups(FirstHandleCallResult).get();
+ // This generates
+ // sycl_handle_special_kernel_parameters<kernel-name-type>(kernelFunc.sout).
+ // (X, Y)
ExprResult Result = SemaRef.BuildCallExpr(
SemaRef.getCurScope(), BaseForSubsCall, Loc, SpecialArgs, Loc);
if (Result.isInvalid())
return true;
- SmallVector<Stmt *> Stmts;
// Make sure to push kernel argument processing result first, before the
- // transformed body of skep-attributed function.
+ // body of skep-attributed function.
+ SmallVector<Stmt *> Stmts;
Stmts.push_back(SemaRef.MaybeCreateExprWithCleanups(Result).get());
- Stmts.push_back(TransformedBody);
- OFDBody = CompoundStmt::Create(SemaRef.getASTContext(), Stmts,
- FPOptionsOverride(), Loc, Loc);
+ Stmts.push_back(Body);
+ ModifiedBody = CompoundStmt::Create(SemaRef.getASTContext(), Stmts,
+ FPOptionsOverride(), Loc, Loc);
} else {
- OFDBody = TransformedBody;
+ ModifiedBody = Body;
}
- OFD->setBody(OFDBody);
+ OutlinedFunctionDeclBodyInstantiator OFDBodyInstantiator(SemaRef, ParmMap,
+ FD);
+ Stmt *TransformedBody = OFDBodyInstantiator.TransformStmt(ModifiedBody).get();
+
+ OFD->setBody(TransformedBody);
OFD->setNothrow();
return OFD;
-
}
class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
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 8857518856c57..7b9524045363c 100644
--- a/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp
+++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp
@@ -72,7 +72,11 @@ void skep2<KN<2>>(K<2>);
// CHECK-NEXT: | | | |-CompoundStmt {{.*}}
// CHECK-NEXT: | | | | `-CallExpr {{.*}} '<dependent type>'
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'KT' lvalue ParmVar {{.*}} 'k' 'KT'
-// CHECK-NEXT: | | | `-UnresolvedLookupExpr {{.*}} '<dependent type>' lvalue (ADL) = 'sycl_kernel_launch' {{.*}}
+// 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: | | | `-UnresolvedLookupExpr {{.*}} '<dependent type>' lvalue (ADL) = 'sycl_handle_special_kernel_parameters' {{.*}}
// CHECK-NEXT: | | | `-TemplateArgument type 'KNT':'type-parameter-0-0'
// CHECK-NEXT: | | | `-TemplateTypeParmType {{.*}} 'KNT' dependent depth 0 index 0
// CHECK-NEXT: | | | `-TemplateTypeParm {{.*}} 'KNT'
@@ -131,7 +135,11 @@ 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: | | | `-UnresolvedLookupExpr {{.*}} '<dependent type>' lvalue (ADL) = 'sycl_kernel_launch' {{.*}}
+// 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: | | | `-UnresolvedLookupExpr {{.*}} '<dependent type>' lvalue (ADL) = 'sycl_handle_special_kernel_parameters' {{.*}}
// CHECK-NEXT: | | | `-TemplateArgument type 'KNT':'type-parameter-0-0'
// CHECK-NEXT: | | | `-TemplateTypeParmType {{.*}} 'KNT' dependent depth 0 index 0
// CHECK-NEXT: | | | `-TemplateTypeParm {{.*}} 'KNT'
@@ -417,7 +425,11 @@ 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: | | | | `-UnresolvedMemberExpr {{.*}} '<bound member function type>' lvalue
+// CHECK-NEXT: | | | | |-UnresolvedMemberExpr {{.*}} '<bound member function type>' lvalue
+// CHECK-NEXT: | | | | `-UnresolvedLookupExpr {{.*}} '<dependent type>' lvalue (ADL) = 'sycl_handle_special_kernel_parameters' {{.*}}
+// 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: | | `-CXXMethodDecl {{.*}} used skep9 {{.*}} implicit_instantiation implicit-inline instantiated_from 0x{{.*}}
// CHECK-NEXT: | | |-TemplateArgument type 'KN<9>'
>From 4dd78ee7dd9f58122a6057b84b58438d87c8624c Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Tue, 14 Jul 2026 03:53:38 -0700
Subject: [PATCH 04/11] Fix checks in pch test
---
.../SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-pch.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
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 047c6268cfe6c..b1f954ee529a3 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
@@ -33,11 +33,11 @@ template void pch_test2<KN<2>>();
#--- test.cpp
// expected-error at +3 {{the 'clang::sycl_kernel_entry_point' kernel name argument conflicts with a previous declaration}}
-// expected-note at pch.h:8 {{previous declaration is here}}
+// expected-note at pch.h:11 {{previous declaration is here}}
[[clang::sycl_kernel_entry_point(KN<1>)]]
void test1() {}
// expected-error at +3 {{the 'clang::sycl_kernel_entry_point' kernel name argument conflicts with a previous declaration}}
-// expected-note at pch.h:12 {{previous declaration is here}}
+// expected-note at pch.h:15 {{previous declaration is here}}
[[clang::sycl_kernel_entry_point(KN<2>)]]
void test2() {}
>From c781130dba231b820caa640fcbd0e0cfd54413af Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Tue, 14 Jul 2026 03:54:06 -0700
Subject: [PATCH 05/11] Fix attribute test
---
clang/test/Misc/pragma-attribute-supported-attributes-list.test | 1 +
1 file changed, 1 insertion(+)
diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
index 8bca68e2119e7..f9f3c15886250 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -201,6 +201,7 @@
// CHECK-NEXT: SYCLExternal (SubjectMatchRule_function)
// CHECK-NEXT: SYCLKernelEntryPoint (SubjectMatchRule_function)
// CHECK-NEXT: SYCLSpecialClass (SubjectMatchRule_record)
+// CHECK-NEXT: SYCLSpecialKernelParameter (SubjectMatchRule_record)
// CHECK-NEXT: ScopedLockable (SubjectMatchRule_record)
// CHECK-NEXT: Section (SubjectMatchRule_function, SubjectMatchRule_variable_is_global, SubjectMatchRule_objc_method, SubjectMatchRule_objc_property)
// CHECK-NEXT: SetTypestate (SubjectMatchRule_function_is_member)
>From 5949c0ff4f372723807dc7cb8689389e472d3220 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Tue, 14 Jul 2026 04:30:24 -0700
Subject: [PATCH 06/11] Add attribute documentation
---
clang/include/clang/Basic/Attr.td | 2 +-
clang/include/clang/Basic/AttrDocs.td | 76 +++++++++++++++++++++++++++
2 files changed, 77 insertions(+), 1 deletion(-)
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 7eb04ee1ffd9f..4e214fd263e3f 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1751,7 +1751,7 @@ def SYCLSpecialKernelParameter : InheritableAttr {
let Spellings = [CXX11<"clang", "sycl_special_kernel_parameter">];
let Subjects = SubjectList<[CXXRecord]>;
let LangOpts = [SYCLHost, SYCLDevice];
- let Documentation = [Undocumented];
+ let Documentation = [SYCLSpecialKernelParameterDocs];
}
def C11NoReturn : InheritableAttr {
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index b67a5b076237c..7e109a16d0234 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -938,6 +938,82 @@ This would trigger the following kernel entry point in the AST:
}];
}
+def SYCLSpecialKernelParameterDocs : Documentation {
+ let Category = DocCatDecl;
+ let Content = [{
+The ``sycl_special_kernel_parameter`` attribute is used to mark a class type
+whose instances require special handling when passed as kernel arguments to a
+function declared with the ``sycl_kernel_entry_point`` attribute. This attribute
+is intended for use by SYCL library implementations to annotate types such as
+``sycl::accessor`` and ``sycl::stream`` that are not trivially copyable and
+therefore cannot be copied to device memory as if by ``memcpy()``.
+
+The attribute only appertains to C++ class or struct types and is only valid
+in SYCL host or device compilation modes.
+
+When a parameter of a function declared with the ``sycl_kernel_entry_point``
+attribute contains a subobject whose type is marked with this attribute, the
+compiler performs special handling on both the host and device sides.
+
+**Host side behavior:**
+
+The synthesized call to ``sycl_kernel_launch`` is expected to return a non-void
+callable object when special kernel parameters are present. This callable is
+then invoked with references to the special subobjects (found by traversing the
+kernel function object's members). The return type of that second call must be
+a ``type_list<Ts...>`` specialization whose template arguments describe the
+types of additional kernel parameters needed to initialize the special objects
+on the device. The synthesized host code looks approximately as follows:
+
+.. code-block:: c++
+
+ // sycl_kernel_launch returns a callable when special params are present,
+ // the callable is invoked with references to special subobjects.
+ type_list<buffer_t*, int> result = sycl_kernel_launch<KN>(
+ symbol, std::move(kernelFunc))(kernelFunc.special_member);
+ // Note that result variable is not declared and is shown here for educational
+ // purposes.
+
+**Device side behavior:**
+
+The generated offload kernel entry point includes additional implicit parameters
+whose types are determined by the ``type_list`` template arguments gathered on
+the host side. These additional parameters are initialized via a synthesized
+call to a ``sycl_handle_special_kernel_parameters`` function template, which
+is looked up in the same manner as ``sycl_kernel_launch``. The synthesized
+device code looks approximately as follows:
+
+.. code-block:: c++
+
+ void kernel-entry-point(KernelType kernelFunc, buffer_t* X, int Y) {
+ sycl_handle_special_kernel_parameters<KernelName>(kernelFunc.special_member)
+ (X, Y);
+ kernelFunc();
+ }
+
+Where ``special_member`` is a subobject of the kernel function object whose type
+is marked with ``sycl_special_kernel_parameter``, and ``X`` and ``Y`` are the
+additional kernel parameters needed to initialize the special subobject on the
+device.
+
+The ``sycl_handle_special_kernel_parameters`` function template is expected to
+be provided by the SYCL library implementation. It receives references to the
+special subobjects as arguments and returns a callable object that accepts the
+additional kernel parameters and uses them to initialize the special subobjects
+for device execution.
+
+Use of this attribute might look as follows:
+
+.. code-block:: c++
+
+ struct [[clang::sycl_special_kernel_parameter]] accessor {
+ int *ptr;
+ size_t range;
+ };
+
+ }];
+}
+
def C11NoReturnDocs : Documentation {
let Category = DocCatFunction;
let Content = [{
>From 76f73c7dc58b3b9cbf6b3c18d31cbaa5c60750d0 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Tue, 14 Jul 2026 06:48:30 -0700
Subject: [PATCH 07/11] Tested base special
---
clang/lib/Sema/SemaSYCL.cpp | 1 -
.../ast-dump-sycl-kernel-decomposition.cpp | 147 +++++++++++++-----
.../CodeGenSYCL/kernel-arg-decomposition.cpp | 55 +++++++
3 files changed, 163 insertions(+), 40 deletions(-)
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index ec62104553406..6af4ccd471bbb 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -595,7 +595,6 @@ class KernelSpecialParamsCreator
}
void visitBaseSpecifierPost(CXXBaseSpecifier *BS) {
- // TODO: test bases
if (isSyclSpecialType(BS->getType()))
createResultingArg();
diff --git a/clang/test/ASTSYCL/ast-dump-sycl-kernel-decomposition.cpp b/clang/test/ASTSYCL/ast-dump-sycl-kernel-decomposition.cpp
index 67e7b0266dc32..08d2d9ca5acf3 100644
--- a/clang/test/ASTSYCL/ast-dump-sycl-kernel-decomposition.cpp
+++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-decomposition.cpp
@@ -73,69 +73,138 @@ template <typename KN, typename KT>
// CHECK-NEXT: | | | `-TemplateTypeParmType {{.*}} 'KN' dependent depth 0 index 0
// CHECK-NEXT: | | | `-TemplateTypeParm {{.*}} 'KN'
// CHECK-NEXT: | | `-SYCLKernelEntryPointAttr {{.*}} KN
-// CHECK-NEXT: | `-FunctionDecl {{.*}} used k {{.*}} implicit_instantiation instantiated_from {{.*}}
-// CHECK-NEXT: | |-TemplateArgument type 'KN<0>'
-// CHECK-NEXT: | | `-RecordType {{.*}} 'KN<0>' canonical
-// CHECK-NEXT: | | `-ClassTemplateSpecialization {{.*}} 'KN'
-// CHECK-NEXT: | |-TemplateArgument type '{{.*}}'
-// CHECK-NEXT: | | `-RecordType {{.*}} canonical
-// CHECK-NEXT: | | `-CXXRecord {{.*}}
-// CHECK-NEXT: | |-ParmVarDecl {{.*}} used Kernel {{.*}}
-// CHECK-NEXT: | |-SYCLKernelCallStmt {{.*}}
+// CHECK-NEXT: | |-FunctionDecl {{.*}} used k {{.*}} implicit_instantiation instantiated_from {{.*}}
+// CHECK-NEXT: | | |-TemplateArgument type 'KN<0>'
+// CHECK-NEXT: | | | `-RecordType {{.*}} 'KN<0>' canonical
+// CHECK-NEXT: | | | `-ClassTemplateSpecialization {{.*}} 'KN'
+// CHECK-NEXT: | | |-TemplateArgument type '{{.*}}'
+// CHECK-NEXT: | | | `-RecordType {{.*}} canonical
+// CHECK-NEXT: | | | `-CXXRecord {{.*}}
+// CHECK-NEXT: | | |-ParmVarDecl {{.*}} used Kernel {{.*}}
+// CHECK-NEXT: | | |-SYCLKernelCallStmt {{.*}}
+// CHECK-NEXT: | | | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | | `-CXXOperatorCallExpr {{.*}} 'void' '()'
+// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} 'void (*)() const' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod {{.*}} 'operator()' 'void () const'
+// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'const {{.*}}' lvalue <NoOp>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} lvalue ParmVar {{.*}} 'Kernel' {{.*}}
+// CHECK-NEXT: | | | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | | `-ExprWithCleanups {{.*}} 'type_list<{{.*}}>'
+// CHECK-NEXT: | | | | `-CXXOperatorCallExpr {{.*}} 'type_list<{{.*}}>' '()'
+// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} 'type_list<{{.*}}> (*)(EmptySpecial &) const' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} 'type_list<{{.*}}> (EmptySpecial &) const' lvalue CXXMethod {{.*}} 'operator()' '{{.*}}'
+// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} 'const {{.*}}' lvalue <NoOp>
+// CHECK-NEXT: | | | | | `-MaterializeTemporaryExpr {{.*}} '{{.*}}' lvalue
+// CHECK-NEXT: | | | | | `-CallExpr {{.*}} '{{.*}}'
+// CHECK-NEXT: | | | | | |-ImplicitCastExpr {{.*}} '{{.*}}' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | | | `-DeclRefExpr {{.*}} '{{.*}}' lvalue Function {{.*}} 'sycl_kernel_launch' {{.*}}
+// CHECK-NEXT: | | | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
+// CHECK-NEXT: | | | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi0EE"
+// CHECK-NEXT: | | | | | `-CXXConstructExpr {{.*}} '{{.*}}' 'void ({{.*}} &&) noexcept'
+// CHECK-NEXT: | | | | | `-ImplicitCastExpr {{.*}} '{{.*}}' xvalue <NoOp>
+// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} lvalue ParmVar {{.*}} 'Kernel' {{.*}}
+// CHECK-NEXT: | | | | `-MemberExpr {{.*}} 'EmptySpecial' lvalue .data {{.*}}
+// CHECK-NEXT: | | | | `-MemberExpr {{.*}} 'Wrapper<EmptySpecial>' lvalue . {{.*}}
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} lvalue ParmVar {{.*}} 'Kernel' {{.*}}
+// CHECK-NEXT: | | | `-OutlinedFunctionDecl {{.*}}
+// CHECK-NEXT: | | | |-ImplicitParamDecl {{.*}} implicit used Kernel {{.*}}
+// CHECK-NEXT: | | | |-ImplicitParamDecl {{.*}} implicit used idk 'int'
+// CHECK-NEXT: | | | `-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | |-ExprWithCleanups {{.*}} 'void'
+// CHECK-NEXT: | | | | `-CXXOperatorCallExpr {{.*}} 'void' '()'
+// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} 'void (*)(int) const' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} 'void (int) const' lvalue CXXMethod {{.*}} 'operator()' '{{.*}}'
+// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} 'const {{.*}}' lvalue <NoOp>
+// CHECK-NEXT: | | | | | `-MaterializeTemporaryExpr {{.*}} '{{.*}}' lvalue
+// CHECK-NEXT: | | | | | `-CallExpr {{.*}} '{{.*}}'
+// CHECK-NEXT: | | | | | |-ImplicitCastExpr {{.*}} '{{.*}}' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | | | `-DeclRefExpr {{.*}} '{{.*}}' lvalue Function {{.*}} 'sycl_handle_special_kernel_parameters' {{.*}}
+// CHECK-NEXT: | | | | | `-CXXConstructExpr {{.*}} 'EmptySpecial' 'void (const EmptySpecial &) noexcept'
+// CHECK-NEXT: | | | | | `-ImplicitCastExpr {{.*}} 'const EmptySpecial' lvalue <NoOp>
+// CHECK-NEXT: | | | | | `-MemberExpr {{.*}} 'EmptySpecial' lvalue .data {{.*}}
+// CHECK-NEXT: | | | | | `-MemberExpr {{.*}} 'Wrapper<EmptySpecial>' lvalue . {{.*}}
+// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} lvalue ImplicitParam {{.*}} 'Kernel' {{.*}}
+// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'int' lvalue ImplicitParam {{.*}} 'idk' 'int'
+// CHECK-NEXT: | | | `-CompoundStmt {{.*}}
+// CHECK-NEXT: | | | `-CXXOperatorCallExpr {{.*}} 'void' '()'
+// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)() const' <FunctionToPointerDecay>
+// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod {{.*}} 'operator()' 'void () const'
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const {{.*}}' lvalue <NoOp>
+// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} lvalue ImplicitParam {{.*}} 'Kernel' {{.*}}
+// CHECK-NEXT: | | `-SYCLKernelEntryPointAttr {{.*}} struct KN<0>
+
+// Test that a class inheriting from a sycl_special_kernel_parameter type
+// is properly decomposed via a DerivedToBase cast.
+// The instantiation for KN<1> should produce a SYCLKernelCallStmt where
+// the base class SpecialBase is accessed via DerivedToBase cast and
+// the additional kernel parameter has struct type SpecialArgData.
+// CHECK: | `-FunctionDecl {{.*}} used k {{.*}} implicit_instantiation instantiated_from {{.*}}
+// CHECK-NEXT: | |-TemplateArgument type 'KN<1>'
+// CHECK: | |-SYCLKernelCallStmt {{.*}}
// CHECK-NEXT: | | |-CompoundStmt {{.*}}
// CHECK-NEXT: | | | `-CXXOperatorCallExpr {{.*}} 'void' '()'
-// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)() const' <FunctionToPointerDecay>
-// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod {{.*}} 'operator()' 'void () const'
-// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const {{.*}}' lvalue <NoOp>
-// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} lvalue ParmVar {{.*}} 'Kernel' {{.*}}
-// CHECK-NEXT: | | |-CompoundStmt {{.*}}
+// CHECK: | | |-CompoundStmt {{.*}}
// CHECK-NEXT: | | | `-ExprWithCleanups {{.*}} 'type_list<{{.*}}>'
// CHECK-NEXT: | | | `-CXXOperatorCallExpr {{.*}} 'type_list<{{.*}}>' '()'
-// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'type_list<{{.*}}> (*)(EmptySpecial &) const' <FunctionToPointerDecay>
-// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'type_list<{{.*}}> (EmptySpecial &) const' lvalue CXXMethod {{.*}} 'operator()' '{{.*}}'
-// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const {{.*}}' lvalue <NoOp>
+// CHECK: | | | |-ImplicitCastExpr {{.*}} 'const {{.*}}' lvalue <NoOp>
// CHECK-NEXT: | | | | `-MaterializeTemporaryExpr {{.*}} '{{.*}}' lvalue
// CHECK-NEXT: | | | | `-CallExpr {{.*}} '{{.*}}'
// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} '{{.*}}' <FunctionToPointerDecay>
// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} '{{.*}}' lvalue Function {{.*}} 'sycl_kernel_launch' {{.*}}
// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} 'const char *' <ArrayToPointerDecay>
-// CHECK-NEXT: | | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi0EE"
+// CHECK-NEXT: | | | | | `-StringLiteral {{.*}} 'const char[14]' lvalue "_ZTS2KNILi1EE"
// CHECK-NEXT: | | | | `-CXXConstructExpr {{.*}} '{{.*}}' 'void ({{.*}} &&) noexcept'
// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} '{{.*}}' xvalue <NoOp>
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} lvalue ParmVar {{.*}} 'Kernel' {{.*}}
-// CHECK-NEXT: | | | `-MemberExpr {{.*}} 'EmptySpecial' lvalue .data {{.*}}
-// CHECK-NEXT: | | | `-MemberExpr {{.*}} 'Wrapper<EmptySpecial>' lvalue . {{.*}}
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'SpecialBase' lvalue <DerivedToBase (SpecialBase)>
+// CHECK-NEXT: | | | `-MemberExpr {{.*}} 'DerivedFromSpecial' lvalue . {{.*}}
// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} lvalue ParmVar {{.*}} 'Kernel' {{.*}}
// CHECK-NEXT: | | `-OutlinedFunctionDecl {{.*}}
// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used Kernel {{.*}}
-// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used idk {{.*}}
+// CHECK-NEXT: | | |-ImplicitParamDecl {{.*}} implicit used idk 'SpecialArgData'
// CHECK-NEXT: | | `-CompoundStmt {{.*}}
// CHECK-NEXT: | | |-ExprWithCleanups {{.*}} 'void'
// CHECK-NEXT: | | | `-CXXOperatorCallExpr {{.*}} 'void' '()'
-// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'void (*)(int) const' <FunctionToPointerDecay>
-// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} 'void (int) const' lvalue CXXMethod {{.*}} 'operator()' '{{.*}}'
-// CHECK-NEXT: | | | |-ImplicitCastExpr {{.*}} 'const {{.*}}' lvalue <NoOp>
-// CHECK-NEXT: | | | | `-MaterializeTemporaryExpr {{.*}} '{{.*}}' lvalue
-// CHECK-NEXT: | | | | `-CallExpr {{.*}} '{{.*}}'
-// CHECK-NEXT: | | | | |-ImplicitCastExpr {{.*}} '{{.*}}' <FunctionToPointerDecay>
-// CHECK-NEXT: | | | | | `-DeclRefExpr {{.*}} '{{.*}}' lvalue Function {{.*}} 'sycl_handle_special_kernel_parameters' {{.*}}
-// CHECK-NEXT: | | | | `-CXXConstructExpr {{.*}} 'EmptySpecial' 'void (const EmptySpecial &) noexcept'
-// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'const EmptySpecial' lvalue <NoOp>
-// CHECK-NEXT: | | | | `-MemberExpr {{.*}} 'EmptySpecial' lvalue .data {{.*}}
-// CHECK-NEXT: | | | | `-MemberExpr {{.*}} 'Wrapper<EmptySpecial>' lvalue . {{.*}}
+// CHECK: | | | | `-DeclRefExpr {{.*}} '{{.*}}' lvalue Function {{.*}} 'sycl_handle_special_kernel_parameters' {{.*}}
+// CHECK-NEXT: | | | | `-CXXConstructExpr {{.*}} 'SpecialBase' 'void (const SpecialBase &) noexcept'
+// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'const SpecialBase' lvalue <DerivedToBase (SpecialBase)>
+// CHECK-NEXT: | | | | `-ImplicitCastExpr {{.*}} 'const DerivedFromSpecial' lvalue <NoOp>
+// CHECK-NEXT: | | | | `-MemberExpr {{.*}} 'DerivedFromSpecial' lvalue . {{.*}}
// CHECK-NEXT: | | | | `-DeclRefExpr {{.*}} lvalue ImplicitParam {{.*}} 'Kernel' {{.*}}
-// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
-// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'int' lvalue ImplicitParam {{.*}} 'idk' 'int'
+// CHECK-NEXT: | | | `-CXXConstructExpr {{.*}} 'SpecialArgData' 'void (const SpecialArgData &) noexcept'
+// CHECK-NEXT: | | | `-ImplicitCastExpr {{.*}} 'const SpecialArgData' lvalue <NoOp>
+// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'SpecialArgData' lvalue ImplicitParam {{.*}} 'idk' 'SpecialArgData'
// CHECK-NEXT: | | `-CompoundStmt {{.*}}
// CHECK-NEXT: | | `-CXXOperatorCallExpr {{.*}} 'void' '()'
-// CHECK-NEXT: | | |-ImplicitCastExpr {{.*}} 'void (*)() const' <FunctionToPointerDecay>
-// CHECK-NEXT: | | | `-DeclRefExpr {{.*}} 'void () const' lvalue CXXMethod {{.*}} 'operator()' 'void () const'
-// CHECK-NEXT: | | `-ImplicitCastExpr {{.*}} 'const {{.*}}' lvalue <NoOp>
+// CHECK: | | `-ImplicitCastExpr {{.*}} 'const {{.*}}' lvalue <NoOp>
// CHECK-NEXT: | | `-DeclRefExpr {{.*}} lvalue ImplicitParam {{.*}} 'Kernel' {{.*}}
-// CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} struct KN<0>
+// CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} struct KN<1>
void case1() {
Wrapper<EmptySpecial> KernelArg;
k<KN<0>>([KernelArg](){});
}
-// CHECK: `-FunctionDecl {{.*}} case1 'void ()'
+
+struct SpecialArgData {
+ int *ptr;
+ int size;
+};
+
+struct [[clang::sycl_special_kernel_parameter]] SpecialBase {
+ int data;
+};
+
+
+struct DerivedFromSpecial : SpecialBase {
+ int extra;
+};
+
+auto set_kernel_arg(SpecialBase &a) {
+ return SpecialArgData{};
+}
+
+void case2() {
+ DerivedFromSpecial DFS;
+ k<KN<1>>([DFS](){});
+}
diff --git a/clang/test/CodeGenSYCL/kernel-arg-decomposition.cpp b/clang/test/CodeGenSYCL/kernel-arg-decomposition.cpp
index b7aa38cc9a1cb..b839e84d9461a 100644
--- a/clang/test/CodeGenSYCL/kernel-arg-decomposition.cpp
+++ b/clang/test/CodeGenSYCL/kernel-arg-decomposition.cpp
@@ -50,6 +50,28 @@ void case1() {
kernel_entry_point<KN<0>>([KernelArg](){});
}
+struct SpecialArgData {
+ int *ptr;
+ int size;
+};
+
+struct [[clang::sycl_special_kernel_parameter]] SpecialBase {
+ int data;
+};
+
+struct DerivedFromSpecial : SpecialBase {
+ int extra;
+};
+
+auto set_kernel_arg(SpecialBase &a) {
+ return SpecialArgData{};
+}
+
+void case2() {
+ DerivedFromSpecial DFS;
+ kernel_entry_point<KN<1>>([DFS](){});
+}
+
// CHECK-HOST-LABEL: define internal void @_Z18kernel_entry_pointI2KNILi0EEZ5case1vEUlvE_EvT0_(
// CHECK-HOST-SAME: i32 [[KERNEL_COERCE0:%.*]], ptr [[KERNEL_COERCE1:%.*]])
// CHECK-HOST: [[ENTRY:.*:]]
@@ -94,3 +116,36 @@ void case1() {
// CHECK-DEVICE: call spir_func void @_ZZ5case1vENKUlvE_clEv(ptr addrspace(4) noundef align 8 dereferenceable_or_null(16) [[KERNEL_ASCAST]])
// CHECK-DEVICE: ret void
+// Test that a class inheriting from a sycl_special_kernel_parameter type
+// produces a kernel entry point with a struct-typed additional parameter
+// (SpecialArgData) and accesses the SpecialBase via DerivedToBase cast.
+
+// CHECK-HOST-LABEL: define internal void @_Z18kernel_entry_pointI2KNILi1EEZ5case2vEUlvE_EvT0_(
+// CHECK-HOST-SAME: i64 [[KERNEL_COERCE:%.*]])
+// CHECK-HOST: [[ENTRY:.*:]]
+// CHECK-HOST: call void @_Z18sycl_kernel_launchI2KNILi1EEJZ5case2vEUlvE_EEDaPKcDpT0_(ptr noundef @.str.1, i64 %{{.*}})
+// CHECK-HOST: [[TMP1:%.*]] = getelementptr inbounds nuw %class.anon.0, ptr %Kernel, i32 0, i32 0
+// CHECK-HOST-NEXT: call void @_ZZ18sycl_kernel_launchI2KNILi1EEJZ5case2vEUlvE_EEDaPKcDpT0_ENKUlDpOT_E_clIJR11SpecialBaseEEEDaS9_(ptr noundef nonnull align 1 dereferenceable(1) %ref.tmp, ptr noundef nonnull align 4 dereferenceable(4) [[TMP1]])
+// CHECK-HOST-NEXT: ret void
+
+// CHECK-DEVICE: define spir_kernel void @_ZTS2KNILi1EE(ptr noundef byval(%class.anon.1) align 4 [[KERNEL:%.*]], ptr noundef byval(%struct.SpecialArgData) align 8 [[IDK:%.*]])
+// CHECK-DEVICE: [[ENTRY:.*:]]
+// CHECK-DEVICE: [[REF_TMP:%.*]] = alloca [[CLASS_ANON_2:%.*]], align 1
+// CHECK-DEVICE: [[AGG_TMP:%.*]] = alloca [[STRUCT_SPECIALBASE:%.*]], align 4
+// CHECK-DEVICE: [[AGG_TMP1:%.*]] = alloca [[STRUCT_SPECIALARGDATA:%.*]], align 8
+// CHECK-DEVICE: [[REF_TMP_ASCAST:%.*]] = addrspacecast ptr [[REF_TMP]] to ptr addrspace(4)
+// CHECK-DEVICE: [[AGG_TMP_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP]] to ptr addrspace(4)
+// CHECK-DEVICE: [[AGG_TMP1_ASCAST:%.*]] = addrspacecast ptr [[AGG_TMP1]] to ptr addrspace(4)
+// CHECK-DEVICE: [[KERNEL_ASCAST:%.*]] = addrspacecast ptr [[KERNEL]] to ptr addrspace(4)
+// CHECK-DEVICE: [[IDK_ASCAST:%.*]] = addrspacecast ptr [[IDK]] to ptr addrspace(4)
+// CHECK-DEVICE: [[REF_TMP_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[REF_TMP_ASCAST]] to ptr
+// CHECK-DEVICE: [[TMP0:%.*]] = getelementptr inbounds nuw [[CLASS_ANON_1:%.*]], ptr addrspace(4) [[KERNEL_ASCAST]], i32 0, i32 0
+// CHECK-DEVICE: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 4 [[AGG_TMP_ASCAST]], ptr addrspace(4) align 4 [[TMP0]], i64 4, i1 false)
+// CHECK-DEVICE: [[AGG_TMP_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP_ASCAST]] to ptr
+// CHECK-DEVICE: call spir_func void @_Z37sycl_handle_special_kernel_parametersI2KNILi1EEJ11SpecialBaseEEDaDpT0_(ptr dead_on_unwind writable sret([[CLASS_ANON_2]]) align 1 [[REF_TMP_ASCAST_ASCAST]], ptr noundef byval([[STRUCT_SPECIALBASE]]) align 4 [[AGG_TMP_ASCAST_ASCAST]])
+// CHECK-DEVICE: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 8 [[AGG_TMP1_ASCAST]], ptr addrspace(4) align 8 [[IDK_ASCAST]], i64 16, i1 false)
+// CHECK-DEVICE: [[AGG_TMP1_ASCAST_ASCAST:%.*]] = addrspacecast ptr addrspace(4) [[AGG_TMP1_ASCAST]] to ptr
+// CHECK-DEVICE: call spir_func void @_ZZ37sycl_handle_special_kernel_parametersI2KNILi1EEJ11SpecialBaseEEDaDpT0_ENKUlDpT_E_clIJ14SpecialArgDataEEEDaS6_(ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) [[REF_TMP_ASCAST]], ptr noundef byval([[STRUCT_SPECIALARGDATA]]) align 8 [[AGG_TMP1_ASCAST_ASCAST]])
+// CHECK-DEVICE: call spir_func void @_ZZ5case2vENKUlvE_clEv(ptr addrspace(4) noundef align 4 dereferenceable_or_null(8) [[KERNEL_ASCAST]])
+// CHECK-DEVICE: ret void
+
>From b359f6cbd7a1a99407472b1f31c8e4825aee9717 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Wed, 15 Jul 2026 06:39:15 -0700
Subject: [PATCH 08/11] Verify serialization for decomposition
---
.../ast-dump-sycl-kernel-decomposition.cpp | 16 +++++++++++++++-
1 file changed, 15 insertions(+), 1 deletion(-)
diff --git a/clang/test/ASTSYCL/ast-dump-sycl-kernel-decomposition.cpp b/clang/test/ASTSYCL/ast-dump-sycl-kernel-decomposition.cpp
index 08d2d9ca5acf3..a68c83bc430b3 100644
--- a/clang/test/ASTSYCL/ast-dump-sycl-kernel-decomposition.cpp
+++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-decomposition.cpp
@@ -5,8 +5,22 @@
// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-unknown -fsycl-is-host \
// RUN: -ast-dump %s \
// RUN: | FileCheck %s
+//
+// Tests with serialization:
+// RUN: %clang_cc1 -std=c++17 -triple spirv64-unknown-unknown -fsycl-is-device \
+// RUN: -emit-pch -o %t %s
+// RUN: %clang_cc1 -x c++ -std=c++17 -triple spirv64-unknown-unknown -fsycl-is-device \
+// RUN: -include-pch %t -ast-dump-all /dev/null \
+// RUN: | sed -e "s/ <undeserialized declarations>//" -e "s/ imported//" \
+// RUN: | FileCheck %s
+// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-unknown -fsycl-is-host \
+// RUN: -emit-pch -o %t %s
+// RUN: %clang_cc1 -x c++ -std=c++17 -triple x86_64-unknown-unknown -fsycl-is-host \
+// RUN: -include-pch %t -ast-dump-all /dev/null \
+// RUN: | sed -e "s/ <undeserialized declarations>//" -e "s/ imported//" \
+// RUN: | FileCheck %s
-// Thes test validates the AST body produced for functions declared with the
+// This test validates the AST body produced for functions declared with the
// sycl_kernel_entry_point attribute in case an argument of such function
// contains an object that requires decomposition.
>From 5cc50b54193b3256831c7e774015cc0076a979e0 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Thu, 16 Jul 2026 04:29:37 -0700
Subject: [PATCH 09/11] Diagnostic fixes
---
.../clang/Basic/DiagnosticSemaKinds.td | 12 +-
clang/include/clang/Sema/Sema.h | 20 +++
clang/include/clang/Sema/SemaSYCL.h | 10 +-
clang/lib/Frontend/FrontendActions.cpp | 4 +
clang/lib/Sema/SemaDecl.cpp | 18 +-
clang/lib/Sema/SemaSYCL.cpp | 154 ++++++++++--------
clang/lib/Sema/SemaTemplateInstantiate.cpp | 51 +++++-
7 files changed, 182 insertions(+), 87 deletions(-)
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 86b765fdf1fab..6c62bec1c61be 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -13565,12 +13565,14 @@ def err_sycl_entry_point_deduced_return_type : Error<
def note_sycl_runtime_defect : Note<
"this indicates a problem with the SYCL runtime header files; please consider"
" reporting this to your SYCL runtime provider">;
-def note_sycl_kernel_launch_lookup_here : Note<
- "in implicit call to 'sycl_kernel_launch' with template argument %0 required"
+def note_sycl_kernel_implicit_lookup_here : Note<
+ "in implicit call to '%0' with template argument %1 required"
" here">;
-def note_sycl_kernel_launch_overload_resolution_here : Note<
- "in implicit call to 'sycl_kernel_launch' with template argument %0 and"
- " function arguments %1 required here">;
+def note_sycl_kernel_implicit_overload_resolution_here : Note<
+ "in implicit call to '%0' with template argument %1 and"
+ " function arguments %2 required here">;
+def note_sycl_kernel_implicit_call_here : Note<
+ "in implicit call to %0 with arguments %1 required here">;
def err_sycl_entry_point_device_use : Error<
"function %0 cannot be used in device code because it is declared with the"
" %1 attribute">;
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index f7d0d493e7081..db01594129b09 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -13307,9 +13307,29 @@ class Sema final : public SemaBase {
/// template named 'sycl_kernel_launch'.
SYCLKernelLaunchLookup,
+ /// We are performing name lookup for a function template or variable
+ /// template named 'sycl_handle_special_kernel_parameters'.
+ SYCLSpecialParametersHandlerLookup,
+
/// We are performing overload resolution for a call to a function
/// template or variable template named 'sycl_kernel_launch'.
SYCLKernelLaunchOverloadResolution,
+
+ /// We are performing overload resolution for a call to a function
+ /// template or variable template named
+ /// 'sycl_handle_special_kernel_parameters'.
+ SYCLSpecialParametersOverloadResolution,
+
+ /// We are synthesizing a call to a callable object returned by implicit
+ /// call to 'sycl_kernel_launch' intended for SYCL kernel special
+ /// kernel arguments handling. This happens in host code.
+ SYCLKernelHostSpecialParametersHandlerCall,
+
+ /// We are synthesizing a call to a callable object returned by
+ /// implicit call to 'sycl_handle_special_kernel_parameters' intended for
+ /// SYCL kernel special kernel arguments handling. This happens in device
+ /// code.
+ SYCLKernelDeviceSpecialParametersHandlerCall,
} Kind;
/// Whether we're substituting into constraints.
diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h
index 268f31d8947cb..1dbf0435fe0d2 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 "clang/Sema/Sema.h"
#include "llvm/ADT/DenseSet.h"
namespace clang {
@@ -83,19 +84,20 @@ class SemaSYCL : public SemaBase {
/// 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,
- StringRef FuncName);
+ ExprResult
+ SynthesizeSYCLKernelIdExpr(FunctionDecl *FD, QualType KernelName,
+ Sema::CodeSynthesisContext::SynthesisKind SKind);
/// 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().
+ /// call to SynthesizeSYCLKernelIdExpr().
StmtResult BuildSYCLKernelCallStmt(FunctionDecl *FD, CompoundStmt *Body,
Expr *LaunchIdExpr,
Expr *HandleSpecParamsExpr);
/// Builds an UnresolvedSYCLKernelCallStmt to wrap 'Body'. 'LaunchIdExpr'
/// specifies the lookup result returned by a previous call to
- /// BuildSYCLKernelLaunchIdExpr().
+ /// SynthesizeSYCLKernelIdExpr().
StmtResult BuildUnresolvedSYCLKernelCallStmt(CompoundStmt *Body,
Expr *LaunchIdExpr,
Expr *HandleSpecParamsExpr);
diff --git a/clang/lib/Frontend/FrontendActions.cpp b/clang/lib/Frontend/FrontendActions.cpp
index f3fce25b78a8c..832e9d0838a57 100644
--- a/clang/lib/Frontend/FrontendActions.cpp
+++ b/clang/lib/Frontend/FrontendActions.cpp
@@ -476,8 +476,12 @@ class DefaultTemplateInstCallback : public TemplateInstantiationCallback {
return "PartialOrderingTTP";
case CodeSynthesisContext::SYCLKernelLaunchLookup:
return "SYCLKernelLaunchLookup";
+ case CodeSynthesisContext::SYCLSpecialParametersHandlerLookup:
+ return "SYCLSpecialParametersHandlerLookup";
case CodeSynthesisContext::SYCLKernelLaunchOverloadResolution:
return "SYCLKernelLaunchOverloadResolution";
+ case CodeSynthesisContext::SYCLSpecialParametersOverloadResolution:
+ return "SYCLSpecialParametersOverloadResolution";
}
return "";
}
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 2646afbac18fe..ccca003b7684f 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -16488,8 +16488,9 @@ Decl *Sema::ActOnStartOfFunctionDef(Scope *FnBodyScope, Decl *D,
// unresolved) call expression after the function body has been parsed.
const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>();
if (!SKEPAttr->isInvalidAttr()) {
- ExprResult LaunchIdExpr =
- SYCL().BuildSYCLKernelLaunchIdExpr(FD, SKEPAttr->getKernelName(), "sycl_kernel_launch");
+ ExprResult LaunchIdExpr = SYCL().SynthesizeSYCLKernelIdExpr(
+ FD, SKEPAttr->getKernelName(),
+ CodeSynthesisContext::SYCLKernelLaunchLookup);
// 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
@@ -16498,12 +16499,13 @@ Decl *Sema::ActOnStartOfFunctionDef(Scope *FnBodyScope, Decl *D,
// 'LaunchIDExpr' failed, then 'SYCLKernelLaunchIdExpr' will be assigned
// a null pointer value below; that is expected.
getCurFunction()->SYCLKernelLaunchIdExpr = LaunchIdExpr.get();
- if (!LaunchIdExpr.isInvalid() &&
- !LaunchIdExpr.get()->getType()->isVoidType()) {
- ExprResult HSPSPIdExpr = SYCL().BuildSYCLKernelLaunchIdExpr(
- FD, SKEPAttr->getKernelName(),
- "sycl_handle_special_kernel_parameters");
- getCurFunction()->HandleSYCLSpecialParamsIdExpr = HSPSPIdExpr.get();
+ if (!LaunchIdExpr.isInvalid()) {
+ ExprResult SpecialParamsHandlerIdExpr =
+ SYCL().SynthesizeSYCLKernelIdExpr(
+ FD, SKEPAttr->getKernelName(),
+ CodeSynthesisContext::SYCLSpecialParametersHandlerLookup);
+ getCurFunction()->HandleSYCLSpecialParamsIdExpr =
+ SpecialParamsHandlerIdExpr.get();
}
}
}
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 6af4ccd471bbb..faac2f7a612a3 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -425,14 +425,19 @@ void SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD) {
}
}
-ExprResult SemaSYCL::BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD,
- QualType KNT,
- StringRef FuncName) {
+ExprResult SemaSYCL::SynthesizeSYCLKernelIdExpr(
+ FunctionDecl *FD, QualType KNT,
+ Sema::CodeSynthesisContext::SynthesisKind SKind) {
// 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 "
"match the requested function context");
+ assert(
+ SKind == Sema::CodeSynthesisContext::SYCLKernelLaunchLookup ||
+ SKind == Sema::CodeSynthesisContext::SYCLSpecialParametersHandlerLookup &&
+ "Only SYCL lookup is expected");
+
// An appropriate source location is required to emit diagnostics if
// lookup fails to produce an overload set. The desired location is the
// start of the function body, but that is not yet available since the
@@ -440,6 +445,10 @@ ExprResult SemaSYCL::BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD,
// The general location of the function is used instead.
SourceLocation Loc = FD->getLocation();
+ StringRef FuncName =
+ (SKind == Sema::CodeSynthesisContext::SYCLKernelLaunchLookup)
+ ? "sycl_kernel_launch"
+ : "sycl_handle_special_kernel_parameters";
ASTContext &Ctx = SemaRef.getASTContext();
IdentifierInfo &SYCLKernelLaunchID =
Ctx.Idents.get(FuncName, tok::TokenKind::identifier);
@@ -448,9 +457,8 @@ ExprResult SemaSYCL::BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD,
// a template named 'sycl_kernel_launch'. In the event of an error, this
// ensures an appropriate diagnostic note is issued to explain why the
// lookup was performed.
- // FIXME: Extend diagnostics for handle special parameters function
Sema::CodeSynthesisContext CSC;
- CSC.Kind = Sema::CodeSynthesisContext::SYCLKernelLaunchLookup;
+ CSC.Kind = SKind;
CSC.Entity = FD;
Sema::ScopedCodeSynthesisContext ScopedCSC(SemaRef, CSC);
@@ -676,6 +684,7 @@ StmtResult BuildSYCLKernelLaunchCallStmt(
return StmtError();
}
+ Expr *KernelLaunchRes;
// Establish a code synthesis context for the implicit call to
// 'sycl_kernel_launch'.
{
@@ -690,39 +699,46 @@ StmtResult BuildSYCLKernelLaunchCallStmt(
SemaRef.BuildCallExpr(SemaRef.getCurScope(), IdExpr, Loc, Args, Loc);
if (LaunchResult.isInvalid())
return StmtError();
- Expr *BaseForSubsCall =
+ KernelLaunchRes =
SemaRef.MaybeCreateExprWithCleanups(LaunchResult).get();
- if (!BaseForSubsCall->getType()->isVoidType()) {
- // FIXME: diagnose that sycl_kernel_launch call returned a callable
- // object. Default diagnostic here is very uncldear
- llvm::SmallVector<Expr *, 12> SpecialArgs;
- for (auto Param : FD->parameters()) {
- if (Param->getType()->isRecordType())
- createArgumentsForSpecialTypes(SpecialArgs, Param, Loc,
- SemaRef.SYCL());
- }
- ExprResult Result = SemaRef.BuildCallExpr(
- SemaRef.getCurScope(), BaseForSubsCall, Loc, SpecialArgs, Loc);
- if (Result.isInvalid())
- return StmtError();
-
- // Now gather types for device code generation. Callable object returned
- // by sycl_kernel_launch call returns type_list object whose template
- // arguments describe types of additional kernel arguments required for
- // special objects, i.e. SYCL accessors/samplers/streams etc.
- QualType Ty = Result.get()->getType();
- // FIXME: that also needs to be diagnosed somewhere.
- auto *TST = Ty->getAs<TemplateSpecializationType>();
- if (!TST)
- return StmtError();
- for (auto Arg : TST->template_arguments()) {
- SpecialArgTys.push_back(Arg.getAsType().getCanonicalType());
- }
-
- Stmts.push_back(SemaRef.MaybeCreateExprWithCleanups(Result).get());
- } else {
- Stmts.push_back(BaseForSubsCall);
+ }
+ if (!KernelLaunchRes->getType()->isVoidType()) {
+ llvm::SmallVector<Expr *, 12> SpecialArgs;
+ for (auto Param : FD->parameters()) {
+ if (Param->getType()->isRecordType())
+ createArgumentsForSpecialTypes(SpecialArgs, Param, Loc,
+ SemaRef.SYCL());
}
+ // Establish a code synthesis context for the implicit call to callable
+ // object returned by 'sycl_kernel_launch'.
+ Sema::CodeSynthesisContext CSC;
+ CSC.Kind = Sema::CodeSynthesisContext::
+ SYCLKernelHostSpecialParametersHandlerCall;
+ CSC.Entity = FD;
+ CSC.CallArgs = SpecialArgs.data();
+ CSC.NumCallArgs = SpecialArgs.size();
+ Sema::ScopedCodeSynthesisContext ScopedCSC(SemaRef, CSC);
+
+ ExprResult Result = SemaRef.BuildCallExpr(
+ SemaRef.getCurScope(), KernelLaunchRes, Loc, SpecialArgs, Loc);
+ if (Result.isInvalid())
+ return StmtError();
+
+ // Now gather types for device code generation. Callable object returned
+ // by sycl_kernel_launch call returns type_list object whose template
+ // arguments describe types of additional kernel arguments required for
+ // special objects, i.e. SYCL accessors/samplers/streams etc.
+ QualType Ty = Result.get()->getType();
+ // FIXME: that also needs to be diagnosed somewhere.
+ auto *TST = Ty->getAs<TemplateSpecializationType>();
+ if (!TST)
+ return StmtError();
+ for (auto Arg : TST->template_arguments())
+ SpecialArgTys.push_back(Arg.getAsType().getCanonicalType());
+
+ Stmts.push_back(SemaRef.MaybeCreateExprWithCleanups(Result).get());
+ } else {
+ Stmts.push_back(KernelLaunchRes);
}
}
@@ -812,7 +828,7 @@ BuildSYCLKernelEntryPointOutline(Sema &SemaRef, FunctionDecl *FD,
// sycl_handle_special_kernel_parameters call.
// This is synthesizing the following pseudo-code:
// void kernel-entry-point(lambda-from-f kernelFunc, buffer_t* X, int Y) {
- // sycl_handle_special_kernel_parameters<kernel-name-type>(kernelFunc.sout).
+ // sycl_handle_special_kernel_parameters<kernel-name-type>(kernelFunc.sout)
// (X, Y);
// {
// // This is copied body of the orignal skep-attributed function.
@@ -832,15 +848,6 @@ BuildSYCLKernelEntryPointOutline(Sema &SemaRef, FunctionDecl *FD,
createArgumentsForSpecialTypes(HandleArgs, PVD, Loc, SemaRef.SYCL());
}
- // FIXME add better diagnosing.
- // Sema::CodeSynthesisContext CSC;
- // CSC.Kind =
- // Sema::CodeSynthesisContext::SYCLKernelLaunchOverloadResolution;
- // CSC.Entity = FD;
- // CSC.CallArgs = Args.data();
- // CSC.NumCallArgs = Args.size();
- // Sema::ScopedCodeSynthesisContext ScopedCSC(SemaRef, CSC);
-
// SpecialArgs are additional kernel arguments that are needed to
// initialize special subobjects and they go to the subsequent call.
// These are (X, Y) in the pseudo code above.
@@ -858,6 +865,7 @@ BuildSYCLKernelEntryPointOutline(Sema &SemaRef, FunctionDecl *FD,
SpecialArgs.push_back(Arg.get());
// Also create an implicit parameter for the device function.
+ // FIXME : IDK????
ImplicitParamDecl *IPD = ImplicitParamDecl::Create(
SemaRef.getASTContext(), OFD, SourceLocation(),
&SemaRef.getASTContext().Idents.get("idk"), QT,
@@ -867,24 +875,42 @@ BuildSYCLKernelEntryPointOutline(Sema &SemaRef, FunctionDecl *FD,
ParmMap[VD] = IPD;
}
- // This generates
- // sycl_handle_special_kernel_parameters<kernel-name-type>(kernelFunc.sout)
- ExprResult FirstHandleCallResult = SemaRef.BuildCallExpr(
- SemaRef.getCurScope(), IdExpr, Loc, HandleArgs, Loc);
- if (FirstHandleCallResult.isInvalid())
- return true;
- // FIXME: diagnose that sycl_special_kernel_parameter call returned a
- // callable object. Default diagnostic here is very uncldear
-
- Expr *BaseForSubsCall =
- SemaRef.MaybeCreateExprWithCleanups(FirstHandleCallResult).get();
- // This generates
- // sycl_handle_special_kernel_parameters<kernel-name-type>(kernelFunc.sout).
- // (X, Y)
- ExprResult Result = SemaRef.BuildCallExpr(
- SemaRef.getCurScope(), BaseForSubsCall, Loc, SpecialArgs, Loc);
- if (Result.isInvalid())
- return true;
+ Expr *BaseForSubsCall;
+ {
+ Sema::CodeSynthesisContext CSC;
+ CSC.Kind =
+ Sema::CodeSynthesisContext::SYCLSpecialParametersOverloadResolution;
+ CSC.Entity = FD;
+ CSC.CallArgs = HandleArgs.data();
+ CSC.NumCallArgs = HandleArgs.size();
+ Sema::ScopedCodeSynthesisContext ScopedCSC(SemaRef, CSC);
+ // This generates
+ // sycl_handle_special_kernel_parameters<kernel-name-type>(kernelFunc.sout)
+ ExprResult FirstHandleCallResult = SemaRef.BuildCallExpr(
+ SemaRef.getCurScope(), IdExpr, Loc, HandleArgs, Loc);
+ if (FirstHandleCallResult.isInvalid())
+ return true;
+ BaseForSubsCall =
+ SemaRef.MaybeCreateExprWithCleanups(FirstHandleCallResult).get();
+ }
+
+ ExprResult Result;
+ {
+ Sema::CodeSynthesisContext CSC;
+ CSC.Kind = Sema::CodeSynthesisContext::
+ SYCLKernelDeviceSpecialParametersHandlerCall;
+ CSC.Entity = FD;
+ CSC.CallArgs = SpecialArgs.data();
+ CSC.NumCallArgs = SpecialArgs.size();
+ Sema::ScopedCodeSynthesisContext ScopedCSC(SemaRef, CSC);
+ // This generates
+ // sycl_handle_special_kernel_parameters<kernel-name-type>(kernelFunc.sout)
+ // (X, Y)
+ Result = SemaRef.BuildCallExpr(
+ SemaRef.getCurScope(), BaseForSubsCall, Loc, SpecialArgs, Loc);
+ if (Result.isInvalid())
+ return true;
+ }
// Make sure to push kernel argument processing result first, before the
// body of skep-attributed function.
diff --git a/clang/lib/Sema/SemaTemplateInstantiate.cpp b/clang/lib/Sema/SemaTemplateInstantiate.cpp
index a77ea5fd3dfff..1da5f62a28b7e 100644
--- a/clang/lib/Sema/SemaTemplateInstantiate.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp
@@ -592,7 +592,11 @@ bool Sema::CodeSynthesisContext::isInstantiationRecord() const {
case TypeAliasTemplateInstantiation:
case PartialOrderingTTP:
case SYCLKernelLaunchLookup:
+ case SYCLSpecialParametersHandlerLookup:
case SYCLKernelLaunchOverloadResolution:
+ case SYCLSpecialParametersOverloadResolution:
+ case SYCLKernelHostSpecialParametersHandlerCall:
+ case SYCLKernelDeviceSpecialParametersHandlerCall:
return false;
// This function should never be called when Kind's value is Memoization.
@@ -1266,33 +1270,68 @@ void Sema::PrintInstantiationStack(InstantiationContextDiagFuncRef DiagFunc) {
<< /*isTemplateTemplateParam=*/true
<< Active->InstantiationRange);
break;
- case CodeSynthesisContext::SYCLKernelLaunchLookup: {
+ case CodeSynthesisContext::SYCLKernelLaunchLookup:
+ case CodeSynthesisContext::SYCLSpecialParametersHandlerLookup: {
const auto *SKEPAttr =
Active->Entity->getAttr<SYCLKernelEntryPointAttr>();
assert(SKEPAttr && "Missing sycl_kernel_entry_point attribute");
assert(!SKEPAttr->isInvalidAttr() &&
"sycl_kernel_entry_point attribute is invalid");
DiagFunc(SKEPAttr->getLocation(), PDiag(diag::note_sycl_runtime_defect));
+ std::string Name =
+ (Active->Kind ==
+ CodeSynthesisContext::SYCLKernelLaunchLookup)
+ ? "sycl_kernel_launch"
+ : "sycl_handle_special_kernel_parameters";
DiagFunc(SKEPAttr->getLocation(),
- PDiag(diag::note_sycl_kernel_launch_lookup_here)
- << SKEPAttr->getKernelName());
+ PDiag(diag::note_sycl_kernel_implicit_lookup_here)
+ << Name << SKEPAttr->getKernelName());
break;
}
- case CodeSynthesisContext::SYCLKernelLaunchOverloadResolution: {
+ case CodeSynthesisContext::SYCLKernelLaunchOverloadResolution:
+ case CodeSynthesisContext::SYCLSpecialParametersOverloadResolution: {
const auto *SKEPAttr =
Active->Entity->getAttr<SYCLKernelEntryPointAttr>();
assert(SKEPAttr && "Missing sycl_kernel_entry_point attribute");
assert(!SKEPAttr->isInvalidAttr() &&
"sycl_kernel_entry_point attribute is invalid");
DiagFunc(SKEPAttr->getLocation(), PDiag(diag::note_sycl_runtime_defect));
+ std::string Name =
+ (Active->Kind ==
+ CodeSynthesisContext::SYCLKernelLaunchOverloadResolution)
+ ? "sycl_kernel_launch"
+ : "sycl_handle_special_kernel_parameters";
DiagFunc(SKEPAttr->getLocation(),
- PDiag(diag::note_sycl_kernel_launch_overload_resolution_here)
- << SKEPAttr->getKernelName()
+ PDiag(diag::note_sycl_kernel_implicit_overload_resolution_here)
+ << Name << SKEPAttr->getKernelName()
<< convertCallArgsValueCategoryAndTypeToString(
*this, llvm::ArrayRef(Active->CallArgs,
Active->NumCallArgs)));
break;
}
+ case CodeSynthesisContext::SYCLKernelDeviceSpecialParametersHandlerCall:
+ case CodeSynthesisContext::SYCLKernelHostSpecialParametersHandlerCall: {
+ const auto *SKEPAttr =
+ Active->Entity->getAttr<SYCLKernelEntryPointAttr>();
+ assert(SKEPAttr && "Missing sycl_kernel_entry_point attribute");
+ assert(!SKEPAttr->isInvalidAttr() &&
+ "sycl_kernel_entry_point attribute is invalid");
+ DiagFunc(SKEPAttr->getLocation(), PDiag(diag::note_sycl_runtime_defect));
+ std::string CalleeName = "callable object returned by ";
+ CalleeName +=
+ (Active->Kind ==
+ CodeSynthesisContext::SYCLKernelHostSpecialParametersHandlerCall)
+ ? "'sycl_kernel_launch'"
+ : "'sycl_handle_special_kernel_parameters'";
+ DiagFunc(SKEPAttr->getLocation(),
+ PDiag(diag::note_sycl_kernel_implicit_call_here)
+ << CalleeName
+ << convertCallArgsValueCategoryAndTypeToString(
+ *this, llvm::ArrayRef(Active->CallArgs,
+ Active->NumCallArgs)));
+
+ break;
+ }
}
}
}
>From 54c8baebbafb7f1f2b0d3106cbb39e53a99bc7ec Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Fri, 17 Jul 2026 04:01:31 -0700
Subject: [PATCH 10/11] Add some Sema testing
---
.../SemaSYCL/kernel-param-decomposition.cpp | 209 ++++++++++++++++++
clang/test/SemaSYCL/sycl-kernel-launch.cpp | 36 +++
2 files changed, 245 insertions(+)
create mode 100644 clang/test/SemaSYCL/kernel-param-decomposition.cpp
diff --git a/clang/test/SemaSYCL/kernel-param-decomposition.cpp b/clang/test/SemaSYCL/kernel-param-decomposition.cpp
new file mode 100644
index 0000000000000..9406bd66b32d0
--- /dev/null
+++ b/clang/test/SemaSYCL/kernel-param-decomposition.cpp
@@ -0,0 +1,209 @@
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-host -fcxx-exceptions -verify %s
+// RUN: %clang_cc1 -triple spirv64-unknown-unknown -std=c++17 -fsyntax-only -fsycl-is-device -verify %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only -fsycl-is-host -fcxx-exceptions -verify %s
+// RUN: %clang_cc1 -triple spirv64-unknown-unknown -std=c++20 -fsyntax-only -fsycl-is-device -verify %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++23 -fsyntax-only -fsycl-is-host -fcxx-exceptions -verify %s
+// RUN: %clang_cc1 -triple spirv64-unknown-unknown -std=c++23 -fsyntax-only -fsycl-is-device -verify %s
+
+// Test overload resolution for implicit calls to
+// sycl_handle_special_kernel_parameters<KN>(...) synthesized for functions
+// declared with the sycl_kernel_entry_point attribute when kernel parameters
+// contain types marked with sycl_special_kernel_parameter.
+
+struct [[clang::sycl_special_kernel_parameter]] SpecialType {
+ int data;
+};
+
+template<typename T>
+struct Wrapper {
+ T data;
+};
+
+template<typename T>
+auto set_kernel_arg(const T &t) {
+ return t;
+}
+
+template<typename... Ts>
+struct type_list {};
+
+template <typename KernelName, typename... Ts>
+auto sycl_kernel_launch(const char *, Ts...) {
+ return [&](auto&&... extra_host_args) {
+ return type_list<decltype(set_kernel_arg(extra_host_args))...>{};
+ };
+}
+
+
+////////////////////////////////////////////////////////////////////////////////
+// Valid declarations.
+////////////////////////////////////////////////////////////////////////////////
+
+// sycl_handle_special_kernel_parameters as function template at namespace scope.
+namespace ok1 {
+ template<int> struct KN;
+ template<typename KN, typename... Ts>
+ auto sycl_handle_special_kernel_parameters(Ts...) {
+ return [](auto ...Args){ return; };
+ }
+ template <typename KNT, typename KT>
+ [[clang::sycl_kernel_entry_point(KNT)]] void skep(KT Kernel) {
+ Kernel();
+ }
+ void test() {
+ Wrapper<SpecialType> KernelArg;
+ skep<KN<1>>([KernelArg](){});
+ }
+}
+
+// sycl_handle_special_kernel_parameters with overload set.
+namespace ok2 {
+ template<int> struct KN;
+ template<typename KN>
+ auto sycl_handle_special_kernel_parameters() {
+ return [](auto ...Args){ return; };
+ }
+ template<typename KN, typename... Ts>
+ auto sycl_handle_special_kernel_parameters(Ts...) {
+ return [](auto ...Args){ return; };
+ }
+ template <typename KNT, typename KT>
+ [[clang::sycl_kernel_entry_point(KNT)]] void skep(KT Kernel) {
+ Kernel();
+ }
+ void test() {
+ Wrapper<SpecialType> KernelArg;
+ skep<KN<2>>([KernelArg](){});
+ }
+}
+
+
+////////////////////////////////////////////////////////////////////////////////
+// Invalid declarations.
+////////////////////////////////////////////////////////////////////////////////
+
+// Undeclared sycl_handle_special_kernel_parameters from function template.
+namespace bad1 {
+ template<int> struct BADKN;
+ struct KT {
+ Wrapper<SpecialType> w;
+ void operator()() const;
+ };
+ // expected-error at +5 {{use of undeclared identifier 'sycl_handle_special_kernel_parameters'}}
+ // 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_handle_special_kernel_parameters' with template argument 'bad1::BADKN<1>' and function arguments (lvalue of type 'SpecialType') required here}}
+ template<typename KNT, typename KTT>
+ [[clang::sycl_kernel_entry_point(KNT)]]
+ void skep(KTT Kernel) {
+ Kernel();
+ }
+ // expected-note at +1 {{in instantiation of function template specialization 'bad1::skep<bad1::BADKN<1>, bad1::KT>' requested here}}
+ template void skep<BADKN<1>>(KT);
+}
+
+// No matching function for call to sycl_handle_special_kernel_parameters;
+// not a template.
+namespace bad2 {
+ template<int> struct BADKN;
+ // expected-note at +1 {{declared as a non-template here}}
+ void sycl_handle_special_kernel_parameters(SpecialType);
+ // expected-error at +4 {{'sycl_handle_special_kernel_parameters' 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_handle_special_kernel_parameters' with template argument 'BADKN<2>' required here}}
+ [[clang::sycl_kernel_entry_point(BADKN<2>)]]
+ void skep(Wrapper<SpecialType> k) {
+ (void)k;
+ }
+}
+
+// No matching function for call to sycl_handle_special_kernel_parameters;
+// mismatched function parameter type.
+namespace bad3 {
+ template<int> struct BADKN;
+ struct KT {
+ Wrapper<SpecialType> w;
+ void operator()() const;
+ };
+ // expected-note at +2 {{candidate function template not viable: no known conversion from 'SpecialType' to 'int' for 1st argument}}
+ template<typename KN, typename... Ts>
+ auto sycl_handle_special_kernel_parameters(int) {
+ return [](auto ...Args){ return; };
+ }
+ // expected-error at +5 {{no matching function for call to 'sycl_handle_special_kernel_parameters'}}
+ // 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_handle_special_kernel_parameters' with template argument 'bad3::BADKN<3>' and function arguments (lvalue of type 'SpecialType') required here}}
+ template<typename KNT, typename KTT>
+ [[clang::sycl_kernel_entry_point(KNT)]]
+ void skep(KTT Kernel) {
+ Kernel();
+ }
+ // expected-note at +1 {{in instantiation of function template specialization 'bad3::skep<bad3::BADKN<3>, bad3::KT>' requested here}}
+ template void skep<BADKN<3>>(KT);
+}
+
+// No matching function for call to sycl_handle_special_kernel_parameters;
+// mismatched template parameter kind.
+namespace bad4 {
+ template<int> struct BADKN;
+ // expected-note at +2 {{candidate template ignored: template argument for non-type template parameter must be an expression}}
+ template<int, typename... Ts>
+ auto sycl_handle_special_kernel_parameters(Ts...) {
+ return [](auto ...Args){ return; };
+ }
+ // expected-error at +4 {{no matching function for call to 'sycl_handle_special_kernel_parameters'}}
+ // 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_handle_special_kernel_parameters' with template argument 'BADKN<4>' and function arguments (lvalue of type 'SpecialType') required here}}
+ [[clang::sycl_kernel_entry_point(BADKN<4>)]]
+ void skep(Wrapper<SpecialType> k) {
+ (void)k;
+ }
+}
+
+// sycl_handle_special_kernel_parameters declared after use and not found by ADL.
+namespace bad5 {
+ template<int> struct BADKN;
+ struct KT {
+ Wrapper<SpecialType> w;
+ void operator()() const;
+ };
+ // expected-error at +5 {{call to function 'sycl_handle_special_kernel_parameters' 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 at +2 {{in implicit call to 'sycl_handle_special_kernel_parameters' with template argument 'bad5::BADKN<5>' and function arguments (lvalue of type 'SpecialType') required here}}
+ template<typename KNT, typename KTT>
+ [[clang::sycl_kernel_entry_point(KNT)]]
+ void skep(KTT Kernel) {
+ Kernel();
+ }
+ // expected-note at +2 {{'sycl_handle_special_kernel_parameters' should be declared prior to the call site or in the global namespace}}
+ template<typename KN, typename... Ts>
+ auto sycl_handle_special_kernel_parameters(Ts...) {
+ return [](auto ...Args){ return; };
+ }
+ // expected-note at +1 {{in instantiation of function template specialization 'bad5::skep<bad5::BADKN<5>, bad5::KT>' requested here}}
+ template void skep<BADKN<5>>(KT);
+}
+
+namespace bad6 {
+// Check that if sycl_handle_special_kernel_parameters doesn't return a callable
+// object, and error is emitted.
+template<int> struct BADKN;
+template<typename KN, typename... Ts>
+auto sycl_handle_special_kernel_parameters(Ts...) {
+ // expected-note at +2 {{conversion candidate of type 'void (*)(int)'}}
+ // expected-note at +1 {{candidate function not viable: no known conversion from 'SpecialType' to 'int' for 1st argument}}
+ return [](int){ return; };
+}
+// expected-error at +4 {{no matching function for call to object of type '(lambda at}}
+// 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 callable object returned by 'sycl_handle_special_kernel_parameters' with arguments (lvalue of type 'SpecialType') required here}}
+template <typename KNT, typename KT>
+[[clang::sycl_kernel_entry_point(KNT)]] void skep(KT Kernel) {
+ Kernel();
+}
+void test() {
+ Wrapper<SpecialType> KernelArg;
+ // expected-note at +1 {{in instantiation of function template specialization 'bad6::skep<bad6::BADKN<6>, (lambda at}}
+ skep<BADKN<6>>([KernelArg](){});
+}
+}
+
diff --git a/clang/test/SemaSYCL/sycl-kernel-launch.cpp b/clang/test/SemaSYCL/sycl-kernel-launch.cpp
index 9734fd9cd99b2..4b689c20ddc73 100644
--- a/clang/test/SemaSYCL/sycl-kernel-launch.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-launch.cpp
@@ -561,3 +561,39 @@ namespace bad18 {
// 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>);
}
+
+namespace bad19 {
+// Check that if sycl_kernel_launch doesn't return a callable
+// object, and error is emitted.
+template<int> struct KN;
+
+struct [[clang::sycl_special_kernel_parameter]] Special {
+ int data;
+};
+
+template<typename T>
+struct Wrapper {
+ T data;
+ int *data1;
+};
+
+template<typename... Ts>
+struct type_list {};
+
+template <typename KernelName, typename... Ts>
+auto sycl_kernel_launch(const char *, Ts...) {
+ return nullptr;
+}
+
+// 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 callable object returned by 'sycl_kernel_launch' with arguments (lvalue of type 'bad19::Special') required here}}
+template <typename KN, typename KT>
+[[clang::sycl_kernel_entry_point(KN)]] void k(KT Kernel) { // expected-error {{called object type 'std::nullptr_t' is not a function or function pointer}}
+ Kernel();
+}
+void case1() {
+ Wrapper<Special> KernelArg;
+// expected-note at +1 {{in instantiation of function template specialization 'bad19::k<BADKN<19>, (lambda at}}
+ k<BADKN<19>>([KernelArg](){});
+}
+}
>From 9c4558c53fb2f6562de7b65cba887f727bd579fd Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Fri, 17 Jul 2026 04:19:38 -0700
Subject: [PATCH 11/11] Maybe fix a fixme and a couple of warnings
---
.../clang/Basic/DiagnosticSemaKinds.td | 3 ++
clang/lib/Frontend/FrontendActions.cpp | 4 +++
clang/lib/Sema/SemaSYCL.cpp | 5 +--
.../SemaSYCL/kernel-param-decomposition.cpp | 31 +++++++++++++++++++
4 files changed, 41 insertions(+), 2 deletions(-)
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 03dee1bab20ad..f4f86677e5230 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -13570,6 +13570,9 @@ def note_sycl_kernel_implicit_overload_resolution_here : Note<
" function arguments %2 required here">;
def note_sycl_kernel_implicit_call_here : Note<
"in implicit call to %0 with arguments %1 required here">;
+def err_sycl_kernel_launch_not_type_list : Error<
+ "the callable returned by 'sycl_kernel_launch' must return a class template"
+ " specialization; %0 is not a class template specialization">;
def err_sycl_entry_point_device_use : Error<
"function %0 cannot be used in device code because it is declared with the"
" %1 attribute">;
diff --git a/clang/lib/Frontend/FrontendActions.cpp b/clang/lib/Frontend/FrontendActions.cpp
index 4f8f369ab1b08..e3e57dcc175b8 100644
--- a/clang/lib/Frontend/FrontendActions.cpp
+++ b/clang/lib/Frontend/FrontendActions.cpp
@@ -482,6 +482,10 @@ class DefaultTemplateInstCallback : public TemplateInstantiationCallback {
return "SYCLKernelLaunchOverloadResolution";
case CodeSynthesisContext::SYCLSpecialParametersOverloadResolution:
return "SYCLSpecialParametersOverloadResolution";
+ case CodeSynthesisContext::SYCLKernelHostSpecialParametersHandlerCall:
+ return "SYCLKernelHostSpecialParametersHandlerCall";
+ case CodeSynthesisContext::SYCLKernelDeviceSpecialParametersHandlerCall:
+ return "SYCLKernelDeviceSpecialParametersHandlerCall";
case CodeSynthesisContext::ExpansionStmtInstantiation:
return "ExpansionStmtInstantiation";
}
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index faac2f7a612a3..4acf14b5126a9 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -729,10 +729,11 @@ StmtResult BuildSYCLKernelLaunchCallStmt(
// arguments describe types of additional kernel arguments required for
// special objects, i.e. SYCL accessors/samplers/streams etc.
QualType Ty = Result.get()->getType();
- // FIXME: that also needs to be diagnosed somewhere.
auto *TST = Ty->getAs<TemplateSpecializationType>();
- if (!TST)
+ if (!TST) {
+ SemaRef.Diag(Loc, diag::err_sycl_kernel_launch_not_type_list) << Ty;
return StmtError();
+ }
for (auto Arg : TST->template_arguments())
SpecialArgTys.push_back(Arg.getAsType().getCanonicalType());
diff --git a/clang/test/SemaSYCL/kernel-param-decomposition.cpp b/clang/test/SemaSYCL/kernel-param-decomposition.cpp
index 9406bd66b32d0..1caa1b0670ed7 100644
--- a/clang/test/SemaSYCL/kernel-param-decomposition.cpp
+++ b/clang/test/SemaSYCL/kernel-param-decomposition.cpp
@@ -207,3 +207,34 @@ void test() {
}
}
+// Check that if sycl_kernel_launch callable returns a non-template-specialization
+// type, an error is emitted.
+namespace bad7 {
+ template<int> struct BADKN;
+ struct KT {
+ Wrapper<SpecialType> w;
+ void operator()() const;
+ };
+ template<typename KN, typename... Ts>
+ auto sycl_handle_special_kernel_parameters(Ts...) {
+ return [](auto ...Args){ return; };
+ }
+ template <typename KernelName, typename... Ts>
+ auto sycl_kernel_launch(const char *, Ts...) {
+ // Returns a callable that returns 'int' instead of a type_list<...> specialization.
+ return [&](auto&&... extra_host_args) {
+ return 42;
+ };
+ }
+ // expected-error at +5 {{the callable returned by 'sycl_kernel_launch' must return a class template specialization; 'int' is not a class template specialization}}
+ // 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 callable object returned by 'sycl_kernel_launch' with arguments (lvalue of type 'SpecialType') required here}}
+ template<typename KNT, typename KTT>
+ [[clang::sycl_kernel_entry_point(KNT)]]
+ void skep(KTT Kernel) {
+ Kernel();
+ }
+ // expected-note at +1 {{in instantiation of function template specialization 'bad7::skep<bad7::BADKN<7>, bad7::KT>' requested here}}
+ template void skep<BADKN<7>>(KT);
+}
+
More information about the cfe-commits
mailing list