llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang-modules
Author: Tom Honermann (tahonermann)
<details>
<summary>Changes</summary>
Previously, implicit calls to the `sycl_kernel_launch()` function provided by
the SYCL runtime library passed a kernel name type as the sole template
argument and a name for the generated kernel entry point function as a string
literal for the first function argument. This design had two issues:
- The generated kernel entry point function name was only communicated to the
SYCL runtime library at run-time despite being known at compile-time.
- The interface left little room for future extension.
With this change, a kernel information type that contains both the kernel name
type and the kernel entry point name as members is passed as the sole template
argument and there is no longer a need to pass an implicit function argument.
The kernel information type has the following structure where `kernel-info`,
`kernel-name-type`, and `kernel-entry-point-name` are exposition only names.
```
struct kernel-info {
using kernel_name = kernel-name-type;
static constexpr const char kernel_entry_point_name[] =
"kernel-entry-point-name";
};
```
The name of the kernel information type is formally unspecified and SYCL
runtime authors should not write code that depends on it. Name mangling for
these types is addressed by defining each such type as an explicit
specialization of an implicitly declared class template named with the reserved
identifier `__sycl_kernel_info`. The names generated for kernel entry point
functions remain unspecified.
The kernel information type reserves room for future extension by allowing for
additional compile-time information to be made available as additional members.
For example, SYCL 2020 allows SYCL kernels to be defined with the
`[[sycl::reqd_work_group_size(N)]]` attribute where `N` specifies a semantic
constraint on how the SYCL runtime launches the kernel. Such constraints could
be communicated by, for example, adding a `reqd_work_group_size` data member to
the kernel information type without affecting the function call interface..
---
Patch is 119.20 KiB, truncated to 20.00 KiB below, full version:
https://github.com/llvm/llvm-project/pull/208958.diff
32 Files Affected:
- (modified) clang/include/clang/AST/ASTNodeTraverser.h (+3-1)
- (modified) clang/include/clang/AST/RecursiveASTVisitor.h (+1)
- (modified) clang/include/clang/AST/StmtSYCL.h (+19-5)
- (modified) clang/include/clang/Basic/AttrDocs.td (+23-9)
- (modified) clang/include/clang/Sema/ScopeInfo.h (+10-3)
- (modified) clang/include/clang/Sema/SemaSYCL.h (+74-25)
- (modified) clang/lib/Sema/SemaDecl.cpp (+25-33)
- (modified) clang/lib/Sema/SemaSYCL.cpp (+285-42)
- (modified) clang/lib/Sema/SemaTemplateInstantiate.cpp (+4-2)
- (modified) clang/lib/Sema/TreeTransform.h (+11-2)
- (modified) clang/lib/Serialization/ASTReaderStmt.cpp (+1)
- (modified) clang/lib/Serialization/ASTWriterStmt.cpp (+1)
- (modified) clang/test/ASTSYCL/ast-dump-sycl-kernel-call-stmt.cpp (+54-46)
- (modified) clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp (+3-3)
- (modified) clang/test/ASTSYCL/ast-print-sycl-kernel-call.cpp (+2-2)
- (modified) clang/test/CodeGenSYCL/function-attrs.cpp (+1-1)
- (modified) clang/test/CodeGenSYCL/kernel-caller-entry-point.cpp (+20-37)
- (modified) clang/test/CodeGenSYCL/sycl-kernel-entry-point-exceptions.cpp
(+12-12)
- (modified) clang/test/CodeGenSYCL/sycl-module-id.cpp (+2-2)
- (modified) clang/test/CodeGenSYCL/unique_stable_name_windows_diff.cpp (+2-2)
- (modified) clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp
(+2-2)
- (modified)
clang/test/SemaSYCL/sycl-kernel-entry-point-attr-device-odr-use.cpp (+2-2)
- (modified) clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
(+2-2)
- (modified)
clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-module.cpp (+2-2)
- (modified)
clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name-pch.cpp (+2-2)
- (modified) clang/test/SemaSYCL/sycl-kernel-entry-point-attr-kernel-name.cpp
(+2-2)
- (modified) clang/test/SemaSYCL/sycl-kernel-entry-point-attr-sfinae.cpp (+2-2)
- (modified) clang/test/SemaSYCL/sycl-kernel-entry-point-attr-this.cpp (+2-2)
- (modified) clang/test/SemaSYCL/sycl-kernel-launch-ms-compat.cpp (+6-6)
- (modified) clang/test/SemaSYCL/sycl-kernel-launch.cpp (+102-102)
- (modified) clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp (+2-2)
- (modified) libsycl/include/sycl/__impl/queue.hpp (+8-9)
``````````diff
diff --git a/clang/include/clang/AST/ASTNodeTraverser.h
b/clang/include/clang/AST/ASTNodeTraverser.h
index d184e03355077..061fd8d87199e 100644
--- a/clang/include/clang/AST/ASTNodeTraverser.h
+++ b/clang/include/clang/AST/ASTNodeTraverser.h
@@ -860,8 +860,10 @@ class ASTNodeTraverser
void
VisitUnresolvedSYCLKernelCallStmt(const UnresolvedSYCLKernelCallStmt *Node) {
Visit(Node->getOriginalStmt());
- if (Traversal != TK_IgnoreUnlessSpelledInSource)
+ if (Traversal != TK_IgnoreUnlessSpelledInSource) {
+ Visit(Node->getKernelInfoType());
Visit(Node->getKernelLaunchIdExpr());
+ }
}
void VisitOMPExecutableDirective(const OMPExecutableDirective *Node) {
diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h
b/clang/include/clang/AST/RecursiveASTVisitor.h
index b77443f1fa0ab..cc62f4cd1fd17 100644
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@@ -3029,6 +3029,7 @@ DEF_TRAVERSE_STMT(SYCLUniqueStableNameExpr, {
DEF_TRAVERSE_STMT(UnresolvedSYCLKernelCallStmt, {
if (getDerived().shouldVisitImplicitCode()) {
TRY_TO(TraverseStmt(S->getOriginalStmt()));
+ TRY_TO(TraverseType(S->getKernelInfoType()));
TRY_TO(TraverseStmt(S->getKernelLaunchIdExpr()));
ShouldVisitChildren = false;
}
diff --git a/clang/include/clang/AST/StmtSYCL.h
b/clang/include/clang/AST/StmtSYCL.h
index 79ac88532e143..75acf1bb9c7f6 100644
--- a/clang/include/clang/AST/StmtSYCL.h
+++ b/clang/include/clang/AST/StmtSYCL.h
@@ -103,27 +103,39 @@ class UnresolvedSYCLKernelCallStmt : public Stmt {
private:
Stmt *OriginalStmt = nullptr;
+
+ // The kernel information type to use in synthesized calls to the SYCL
+ // runtime library. This type is a specialization of the class template
+ // returned by SemaSYCL::GetSYCLKernelInfoClassTemplate() specialized for
+ // the kernel name type. This will be a dependent type if the kernel name
+ // type is dependent.
+ QualType KernelInfoType;
+
// KernelLaunchIdExpr stores an UnresolvedLookupExpr or UnresolvedMemberExpr
// corresponding to the SYCL kernel launch function for which a call
// will be synthesized during template instantiation.
Expr *KernelLaunchIdExpr = nullptr;
- UnresolvedSYCLKernelCallStmt(CompoundStmt *CS, Expr *IdExpr)
+ UnresolvedSYCLKernelCallStmt(CompoundStmt *CS, QualType InfoType,
+ Expr *IdExpr)
: Stmt(UnresolvedSYCLKernelCallStmtClass), OriginalStmt(CS),
- KernelLaunchIdExpr(IdExpr) {}
+ KernelInfoType(InfoType), KernelLaunchIdExpr(IdExpr) {}
void setOriginalStmt(CompoundStmt *CS) { OriginalStmt = CS; }
+ void setKernelInfoType(QualType InfoType) { KernelInfoType = InfoType; }
+
void setKernelLaunchIdExpr(Expr *IdExpr) { KernelLaunchIdExpr = IdExpr; }
public:
static UnresolvedSYCLKernelCallStmt *Create(const ASTContext &C,
- CompoundStmt *CS, Expr *IdExpr) {
- return new (C) UnresolvedSYCLKernelCallStmt(CS, IdExpr);
+ CompoundStmt *CS,
+ QualType InfoType, Expr *IdExpr)
{
+ return new (C) UnresolvedSYCLKernelCallStmt(CS, InfoType, IdExpr);
}
static UnresolvedSYCLKernelCallStmt *CreateEmpty(const ASTContext &C) {
- return new (C) UnresolvedSYCLKernelCallStmt(nullptr, nullptr);
+ return new (C) UnresolvedSYCLKernelCallStmt(nullptr, QualType(), nullptr);
}
CompoundStmt *getOriginalStmt() { return cast<CompoundStmt>(OriginalStmt); }
@@ -131,6 +143,8 @@ class UnresolvedSYCLKernelCallStmt : public Stmt {
return cast<CompoundStmt>(OriginalStmt);
}
+ QualType getKernelInfoType() const { return KernelInfoType; }
+
Expr *getKernelLaunchIdExpr() { return KernelLaunchIdExpr; }
const Expr *getKernelLaunchIdExpr() const { return KernelLaunchIdExpr; }
diff --git a/clang/include/clang/Basic/AttrDocs.td
b/clang/include/clang/Basic/AttrDocs.td
index 05e4cb0870652..1a2d9244e064d 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -701,13 +701,23 @@ follows.
namespace sycl {
class handler {
- template<typename KernelName, typename... Ts>
- void sycl_kernel_launch(const char* kernelSymbol, Ts&&... kernelArgs) {
+ template<typename KernelInfo, typename... Ts>
+ void sycl_kernel_launch(Ts&&... kernelArgs) {
// This code will run on the host and is responsible for calling
functions
// appropriate for the desired offload backend (OpenCL, CUDA, HIP,
// Level Zero, etc...) to copy the kernel arguments denoted by kernelArgs
// to a device and to schedule an invocation of the offload kernel entry
- // point denoted by kernelSymbol with the copied arguments.
+ // point with the copied arguments.
+ //
+ // The KernelInfo template parameter denotes an unspecified type which
+ // contains metadata for the kernel. It has the following members:
+ // - kernel_name: a type alias for the kernel name type. This corresponds
+ // to the KernelName template parameter of SYCL kernel invocation
+ // functions.
+ // - kernel_entry_point_name: a static constexpr array of char containing
+ // the name of the generated offload kernel entry point function. The
+ // array is suitable for use in constant expressions, including
non-type
+ // template arguments.
}
template<typename KernelName, typename KernelType>
@@ -823,12 +833,16 @@ There are a few items worthy of note:
copyability and kernel parameter requirements.
Within ``single_task()``, the call to ``kernel_entry_point()`` is effectively
-replaced with a synthesized call to a ''sycl_kernel_launch`` template that
+replaced with a synthesized call to a ``sycl_kernel_launch`` template that
looks approximately as follows.
.. code-block:: c++
- sycl_kernel_launch<KN>("sycl-kernel-caller-for-KN", kernelFunc);
+ struct kernel-info {
+ using kernel_name = KN;
+ static constexpr const char kernel_entry_point_name[] =
"sycl-kernel-caller-for-KN";
+ };
+ sycl_kernel_launch<KN>(kernelFunc);
There are a few items worthy of note:
@@ -844,10 +858,10 @@ There are a few items worthy of note:
#. The ``sycl_kernel_launch`` template is expected to be provided by the SYCL
library implementation. It is responsible for copying the kernel arguments
to device memory and for scheduling execution of the generated offload
- kernel entry point identified by the symbol name passed as the first
- function argument. ``sycl-kernel-caller-for-KN`` is substituted above for
- the actual symbol name that would be generated for the offload kernel entry
- point.
+ kernel entry point identified by the ``kernel_entry_point_name`` member of
+ the first template parameter. ``sycl-kernel-caller-for-KN`` is substituted
+ above for the actual symbol name that would be generated for the offload
+ kernel entry point.
It is not necessary for a function declared with the
``sycl_kernel_entry_point``
attribute to be called for the offload kernel entry point to be emitted. For
diff --git a/clang/include/clang/Sema/ScopeInfo.h
b/clang/include/clang/Sema/ScopeInfo.h
index 7e4d3f2e0d1cb..fb92ee0d2ef49 100644
--- a/clang/include/clang/Sema/ScopeInfo.h
+++ b/clang/include/clang/Sema/ScopeInfo.h
@@ -23,6 +23,7 @@
#include "clang/Basic/SourceLocation.h"
#include "clang/Sema/CleanupInfo.h"
#include "clang/Sema/DeclSpec.h"
+#include "clang/Sema/SemaSYCL.h"
#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/DenseMapInfo.h"
#include "llvm/ADT/MapVector.h"
@@ -245,9 +246,15 @@ class FunctionScopeInfo {
/// The set of GNU address of label extension "&&label".
llvm::SmallVector<AddrLabelExpr *, 4> AddrLabels;
- /// An unresolved identifier lookup expression for an implicit call
- /// to a SYCL kernel launch function in a dependent context.
- Expr *SYCLKernelLaunchIdExpr = nullptr;
+ // Implicit calls to the SYCL runtime library are synthesized for functions
+ // declared with the sycl_kernel_entry_point attribute. These calls require
+ // types and lookup results for SYCL runtime library declared functions that
+ // are produced at the start of a function definition. SYCLKernelASTFragments
+ // caches these artifacts for later use to construct a SYCLKernelCallStmt or
+ // UnresolvedSYCLKernelCallStmt. If production of these artifacts fails, a
+ // diagnostic will have been issued and a disengaged optional stored.
+ std::optional<SemaSYCL::SYCLKernelCallStmtASTFragments>
+ SYCLKernelASTFragments;
public:
/// Represents a simple identification of a weak object.
diff --git a/clang/include/clang/Sema/SemaSYCL.h
b/clang/include/clang/Sema/SemaSYCL.h
index 4980aa44c3012..fff3c65983365 100644
--- a/clang/include/clang/Sema/SemaSYCL.h
+++ b/clang/include/clang/Sema/SemaSYCL.h
@@ -18,6 +18,7 @@
#include "clang/Basic/SourceLocation.h"
#include "clang/Sema/Ownership.h"
#include "clang/Sema/SemaBase.h"
+#include "llvm/ADT/DenseMap.h"
#include "llvm/ADT/DenseSet.h"
namespace clang {
@@ -25,6 +26,16 @@ class Decl;
class ParsedAttr;
class SemaSYCL : public SemaBase {
+ /// SYCLKernelInfoClassTemplate caches the class template declaration built
+ /// by GetSYCLKernelInfoClassTemplate().
+ ClassTemplateDecl *SYCLKernelInfoClassTemplate = nullptr;
+
+ /// SYCLKernelInfoClassTemplateSpecializations maps SYCL kernel name types
+ /// to their corresponding synthesized explicit specialization declarations
+ /// of the SYCL kernel info class template.
+ llvm::DenseMap<CanQualType, ClassTemplateSpecializationDecl *>
+ SYCLKernelInfoClassTemplateSpecializations;
+
public:
SemaSYCL(Sema &S);
@@ -71,31 +82,69 @@ class SemaSYCL : public SemaBase {
void CheckSYCLExternalFunctionDecl(FunctionDecl *FD);
void CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD);
- /// Builds an expression for the lookup of a 'sycl_kernel_launch' template
- /// with 'KernelName' as an explicit template argument. Lookup is performed
- /// as if from the first statement of the body of 'FD' and thus requires
- /// searching the scopes that exist at parse time. This function therefore
- /// requires the current semantic context to be the definition of 'FD'. In a
- /// dependent context, the returned expression will be an
UnresolvedLookupExpr
- /// or an UnresolvedMemberExpr. In a non-dependent context, the returned
- /// expression will be a DeclRefExpr or MemberExpr. If lookup fails, a null
- /// error result is returned. The resulting expression is intended to be
- /// passed as the 'LaunchIdExpr' argument in a call to either
- /// BuildSYCLKernelCallStmt() or BuildUnresolvedSYCLKernelCallStmt() after
- /// the function body has been parsed.
- ExprResult BuildSYCLKernelLaunchIdExpr(FunctionDecl *FD, QualType
KernelName);
-
- /// Builds a SYCLKernelCallStmt to wrap 'Body' and to be used as the body of
- /// 'FD'. 'LaunchIdExpr' specifies the lookup result returned by a previous
- /// call to BuildSYCLKernelLaunchIdExpr().
- StmtResult BuildSYCLKernelCallStmt(FunctionDecl *FD, CompoundStmt *Body,
- Expr *LaunchIdExpr);
-
- /// Builds an UnresolvedSYCLKernelCallStmt to wrap 'Body'. 'LaunchIdExpr'
- /// specifies the lookup result returned by a previous call to
- /// BuildSYCLKernelLaunchIdExpr().
- StmtResult BuildUnresolvedSYCLKernelCallStmt(CompoundStmt *Body,
- Expr *LaunchIdExpr);
+ /// GetSYCLKernelInfoClassTemplate builds a class template used to pass SYCL
+ /// kernel information in synthesized calls to SYCL runtime library
functions.
+ /// The class template has an unspecified name, is an incomplete class
+ /// (explicit specializations are synthesized as needed), built on first
call,
+ /// and cached for return in subsequent calls.
+ ClassTemplateDecl *GetSYCLKernelInfoClassTemplate();
+
+ /// GetSYCLKernelInfoClassSpecializationType returns a type for the class
+ /// template returned by GetSYCLKernelInfoClassTemplate() specialized for
+ /// 'KNT'. 'KNT' may be a dependent type. If 'KNT' is not a dependent type,
+ /// an explicit specialization declaration for the specialization is
+ /// synthesized and a canonical type returned. A definition for the explicit
+ /// specialization is synthesized when SYCLKernelCallStmt is constructed for
+ /// 'KNT'.
+ QualType GetSYCLKernelInfoClassSpecializationType(QualType KNT);
+
+ /// SYCLKernelCallStmtASTFragments holds portions of the AST used for lookup
+ /// of SYCL runtime library functions. These are constructed at the beginning
+ /// of a function definition and later used in synthesized calls to the
+ /// SYCL runtime library. In a dependent context, these are used to construct
+ /// an UnresolvedSYCLKernelCallStmt. During template instantiation, or in a
+ /// non-dependent context, these are used to construct a SYCLKernelCallStmt.
+ ///
+ /// KernelInfoType is the type to be passed as the explicit template argument
+ /// to SYCL runtime libraries.
+ ///
+ /// KernelLaunchIdExpr is an UnresolvedLookupExpr or UnresolvedMemberExpr
+ /// for the SYCL kernel launch function, with KernelInfoType used as the
+ /// explicit template argument.
+ struct SYCLKernelCallStmtASTFragments {
+ QualType KernelInfoType;
+ Expr *KernelLaunchIdExpr;
+ };
+
+ /// BuildSYCLKernelCallStmtASTFragments builds portions of the AST needed
+ /// to construct SYCLKernelCallStmt or UnresolvedSYCLKernelCallStmt that
+ /// must be constructed early in the definition of a function before the
+ /// body of the function has been parsed.
+ ///
+ /// 'FD' must be a function declared with a valid sycl_kernel_entry_point
+ /// attribute and the current Sema context must match 'FD'.
+ ///
+ /// If construction of the AST fragments fails, diagnostics are issued and
+ /// a disengaged optional value is returned.
+ std::optional<SYCLKernelCallStmtASTFragments>
+ BuildSYCLKernelCallStmtASTFragments(FunctionDecl *FD);
+
+ /// BuildSYCLKernelCallStmt constructs a SYCLKernelCallStmt or an
+ /// UnresolvedSYCLKernelCallStmt depending on whether 'FD' is a templated
+ /// function.
+ ///
+ /// 'FD' must be a function declared with a valid sycl_kernel_entry_point
+ /// attribute, must not have had its body assigned yet, and the current
+ /// Sema context must match 'FD'.
+ ///
+ /// 'Body' is the parsed compound statement body of 'FD' to be wrapped by
+ /// the new SYCLKernelCallStmt or UnresolvedSYCLKernelCallStmt.
+ ///
+ /// 'ASTFragments' contains the AST fragments previously returned by a call
+ /// to BuildSYCLKernelCallStmtASTFragments() for 'FD'.
+ StmtResult
+ BuildSYCLKernelCallStmt(FunctionDecl *FD, CompoundStmt *Body,
+ const SYCLKernelCallStmtASTFragments &ASTFragments);
};
} // namespace clang
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index cb876cf8dd936..1d5ccf1c793eb 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -16513,27 +16513,23 @@ Decl *Sema::ActOnStartOfFunctionDef(Scope
*FnBodyScope, Decl *D,
if (!FD->isInvalidDecl() && FD->hasAttr<SYCLKernelEntryPointAttr>() &&
FnBodyScope) {
- // An implicit call expression is synthesized for functions declared with
- // the sycl_kernel_entry_point attribute. The call may resolve to a
- // function template, a member function template, or a call operator
- // of a variable template depending on the results of unqualified lookup
- // for 'sycl_kernel_launch' from the beginning of the function body.
- // Performing that lookup requires the stack of parsing scopes active
- // when the definition is parsed and is thus done here; the result is
- // cached in FunctionScopeInfo and used to synthesize the (possibly
- // unresolved) call expression after the function body has been parsed.
+ // A function declared with the sycl_kernel_entry_point attribute has its
+ // body wrapped in a SYCLKernelCallStmt or UnresolvedSYCLKernelCallStmt.
+ // Those statements cannot be constructed until the function body has been
+ // parsed, but portions of the AST needed for construction have to be
+ // constructed early, before the function body has been parsed. Those are
+ // collected here.
const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>();
if (!SKEPAttr->isInvalidAttr()) {
- ExprResult LaunchIdExpr =
- SYCL().BuildSYCLKernelLaunchIdExpr(FD, SKEPAttr->getKernelName());
- // Do not mark 'FD' as invalid if construction of `LaunchIDExpr` produces
- // an invalid result. Name lookup failure for 'sycl_kernel_launch' is
- // treated as an error in the definition of 'FD'; treating it as an error
- // of the declaration would affect overload resolution which would
- // potentially result in additional errors. If construction of
- // 'LaunchIDExpr' failed, then 'SYCLKernelLaunchIdExpr' will be assigned
- // a null pointer value below; that is expected.
- getCurFunction()->SYCLKernelLaunchIdExpr = LaunchIdExpr.get();
+ // If construction of the AST fragments fails, SYCLKernelASTFragments
+ // will be assigned a disengaged value. This is expected and checked
+ // after the function body is parsed. Failure to construct the AST
+ // fragments is treated as an error in the definition of 'FD'; 'FD'
+ // is not marked as invalid since treating such failures as an error
+ // in the declaration would affect overload resolution which would
+ // potentially result in additional errors.
+ getCurFunction()->SYCLKernelASTFragments =
+ SYCL().BuildSYCLKernelCallStmtASTFragments(FD);
}
}
@@ -16738,30 +16734,26 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt
*Body, bool IsInstantiation,
SKEPAttr->setInvalidAttr();
}
- // Build an unresolved SYCL kernel call statement for a function template,
- // validate that a SYCL kernel call statement was instantiated for an
- // (implicit or explicit) instantiation of a function template, or
otherwise
- // build a (resolved) SYCL kernel call statement for a non-templated
- // function or an explicit specialization.
+ // Build a SYCL kernel call statement to wrap the function body.
if (Body && !SKEPAttr->isInvalidAttr()) {
StmtResult SR;
if (FD->isTemplateInstantiation()) {
// The function body should already be a SYCLKernelCallStmt in this
// case, but might not be if there were previous errors.
SR = Body;
- } else if (!getCurFunction()->SYCLKernelLaunchIdExpr) {
- // If name lookup for a template named sycl_kernel_launch failed
- // earlier, don't try to build a SYCL kernel call statement as that
- // would cause additional errors to be issued; just proceed with the
- // original function body.
+ } else if (!getCurFunction()->SYCLKernelASTFragments) {
+ // Construction of the AST fragments failed earlier and diagnostics
+ // will have already been issued. Don't try to build a SYCL kernel
+ // call statement as that would cause additional errors to be issued;
+ // proceed with the original function body.
SR = Body;
- } else if (FD->isTemplated()) {
- SR = SYCL().BuildUnresolvedSYCLKernelCallStmt(
- cast<CompoundStmt>(Body),
getCurFunction()->SYCLKernelLaunchIdExpr);
} else {
+ // Construct the SYCL kernel call statement. This will build an
+ // UnresolvedSYCLKernelCallStmt if FD is templated and a
+ // SYCLKernelCallStmt otherwise.
SR = SYCL().BuildSYCLKernelCallStmt(
FD, cast<CompoundStmt>(Body),
- getCurFunction()->SYCLKernelLaunchIdExpr);
+ *getCurFunction()->SYCLKernelASTFragments);
}
// If construction of the replacement body fails, just continue with the
// original function body. An early error return here is not valid; the
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index b942f19761f40..4de94621a5f19 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -425,8 +425,204 @@ void
SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD) {
}
}
-ExprResult SemaSYCL::BuildS...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/208958
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
