[clang] 1a59087 - [SYCL] The sycl_kernel_entry_point attribute. (#111389)
via cfe-commits
cfe-commits at lists.llvm.org
Tue Nov 5 08:09:37 PST 2024
Author: Tom Honermann
Date: 2024-11-05T11:09:32-05:00
New Revision: 1a590870b6b3452934ecc245e01957fdab48909c
URL: https://github.com/llvm/llvm-project/commit/1a590870b6b3452934ecc245e01957fdab48909c
DIFF: https://github.com/llvm/llvm-project/commit/1a590870b6b3452934ecc245e01957fdab48909c.diff
LOG: [SYCL] The sycl_kernel_entry_point attribute. (#111389)
The `sycl_kernel_entry_point` attribute is used to declare a function that
defines a pattern for an offload kernel to be emitted. The attribute requires
a single type argument that specifies the type used as a SYCL kernel name as
described in section 5.2, "Naming of kernels", of the SYCL 2020 specification.
Properties of the offload kernel are collected when a function declared with
the `sycl_kernel_entry_point` attribute is parsed or instantiated. These
properties, such as the kernel name type, are stored in the AST context where
they are (or will be) used for diagnostic purposes and to facilitate reflection
to a SYCL run-time library. These properties are not serialized with the AST
but are recreated upon deserialization.
The `sycl_kernel_entry_point` attribute is intended to replace the existing
`sycl_kernel` attribute which is intended to be deprecated in a future change
and removed following an appropriate deprecation period. The new attribute
differs in that it is enabled for both SYCL host and device compilation, may
be used with non-template functions, explicitly indicates the type used as
the kernel name type, and will impact AST generation.
This change adds the basic infrastructure for the new attribute. Future
changes will add diagnostics and new AST support that will be used to drive
generation of the corresponding offload kernel.
Added:
clang/include/clang/AST/SYCLKernelInfo.h
clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
clang/test/SemaSYCL/sycl-kernel-entry-point-attr-ignored.cpp
Modified:
clang/include/clang/AST/ASTContext.h
clang/include/clang/Basic/Attr.td
clang/include/clang/Basic/AttrDocs.td
clang/include/clang/Sema/SemaSYCL.h
clang/lib/AST/ASTContext.cpp
clang/lib/Sema/SemaDecl.cpp
clang/lib/Sema/SemaDeclAttr.cpp
clang/lib/Sema/SemaSYCL.cpp
clang/lib/Serialization/ASTReaderDecl.cpp
clang/test/Misc/pragma-attribute-supported-attributes-list.test
Removed:
################################################################################
diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h
index fb0c051dc19182..1e8101f60b03fb 100644
--- a/clang/include/clang/AST/ASTContext.h
+++ b/clang/include/clang/AST/ASTContext.h
@@ -23,6 +23,7 @@
#include "clang/AST/ExternalASTSource.h"
#include "clang/AST/PrettyPrinter.h"
#include "clang/AST/RawCommentList.h"
+#include "clang/AST/SYCLKernelInfo.h"
#include "clang/AST/TemplateName.h"
#include "clang/Basic/LLVM.h"
#include "clang/Basic/PartialDiagnostic.h"
@@ -1239,6 +1240,11 @@ class ASTContext : public RefCountedBase<ASTContext> {
/// in device compilation.
llvm::DenseSet<const FunctionDecl *> CUDAImplicitHostDeviceFunUsedByDevice;
+ /// Map of SYCL kernels indexed by the unique type used to name the kernel.
+ /// Entries are not serialized but are recreated on deserialization of a
+ /// sycl_kernel_entry_point attributed function declaration.
+ llvm::DenseMap<CanQualType, SYCLKernelInfo> SYCLKernels;
+
/// For capturing lambdas with an explicit object parameter whose type is
/// derived from the lambda type, we need to perform derived-to-base
/// conversion so we can access the captures; the cast paths for that
@@ -3340,6 +3346,14 @@ class ASTContext : public RefCountedBase<ASTContext> {
void getFunctionFeatureMap(llvm::StringMap<bool> &FeatureMap,
GlobalDecl GD) const;
+ /// Generates and stores SYCL kernel metadata for the provided
+ /// SYCL kernel entry point function. The provided function must have
+ /// an attached sycl_kernel_entry_point attribute that specifies a unique
+ /// type for the name of a SYCL kernel. Callers are required to detect
+ /// conflicting SYCL kernel names and issue a diagnostic prior to calling
+ /// this function.
+ void registerSYCLEntryPointFunction(FunctionDecl *FD);
+
//===--------------------------------------------------------------------===//
// Statistics
//===--------------------------------------------------------------------===//
diff --git a/clang/include/clang/AST/SYCLKernelInfo.h b/clang/include/clang/AST/SYCLKernelInfo.h
new file mode 100644
index 00000000000000..55dba1f8e31fd9
--- /dev/null
+++ b/clang/include/clang/AST/SYCLKernelInfo.h
@@ -0,0 +1,41 @@
+//===--- SYCLKernelInfo.h --- Information about SYCL kernels --------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+/// \file
+/// This file declares types used to describe SYCL kernels.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_CLANG_AST_SYCLKERNELINFO_H
+#define LLVM_CLANG_AST_SYCLKERNELINFO_H
+
+#include "clang/AST/Decl.h"
+#include "clang/AST/Type.h"
+
+namespace clang {
+
+class SYCLKernelInfo {
+public:
+ SYCLKernelInfo(CanQualType KernelNameType,
+ const FunctionDecl *KernelEntryPointDecl)
+ : KernelNameType(KernelNameType),
+ KernelEntryPointDecl(KernelEntryPointDecl) {}
+
+ CanQualType getKernelNameType() const { return KernelNameType; }
+
+ const FunctionDecl *getKernelEntryPointDecl() const {
+ return KernelEntryPointDecl;
+ }
+
+private:
+ CanQualType KernelNameType;
+ const FunctionDecl *KernelEntryPointDecl;
+};
+
+} // namespace clang
+
+#endif // LLVM_CLANG_AST_SYCLKERNELINFO_H
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 156fbd1c4442eb..fbad11b376e7e9 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -407,7 +407,8 @@ def MicrosoftExt : LangOpt<"MicrosoftExt">;
def Borland : LangOpt<"Borland">;
def CUDA : LangOpt<"CUDA">;
def HIP : LangOpt<"HIP">;
-def SYCL : LangOpt<"SYCLIsDevice">;
+def SYCLHost : LangOpt<"SYCLIsHost">;
+def SYCLDevice : LangOpt<"SYCLIsDevice">;
def COnly : LangOpt<"", "!LangOpts.CPlusPlus">;
def CPlusPlus : LangOpt<"CPlusPlus">;
def OpenCL : LangOpt<"OpenCL">;
@@ -1493,14 +1494,23 @@ def : MutualExclusions<[CUDAConstant, CUDAShared, HIPManaged]>;
def SYCLKernel : InheritableAttr {
let Spellings = [Clang<"sycl_kernel">];
let Subjects = SubjectList<[FunctionTmpl]>;
- let LangOpts = [SYCL];
+ let LangOpts = [SYCLDevice];
let Documentation = [SYCLKernelDocs];
}
+def SYCLKernelEntryPoint : InheritableAttr {
+ let Spellings = [Clang<"sycl_kernel_entry_point">];
+ let Args = [TypeArgument<"KernelName">];
+ let Subjects = SubjectList<[Function], ErrorDiag>;
+ let TemplateDependent = 1;
+ let LangOpts = [SYCLHost, SYCLDevice];
+ let Documentation = [SYCLKernelEntryPointDocs];
+}
+
def SYCLSpecialClass: InheritableAttr {
let Spellings = [Clang<"sycl_special_class">];
let Subjects = SubjectList<[CXXRecord]>;
- let LangOpts = [SYCL];
+ let LangOpts = [SYCLDevice];
let Documentation = [SYCLSpecialClassDocs];
}
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index b497cce37625c9..ed251b0a74c392 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -455,6 +455,180 @@ The SYCL kernel in the previous code sample meets these expectations.
}];
}
+def SYCLKernelEntryPointDocs : Documentation {
+ let Category = DocCatFunction;
+ let Content = [{
+The ``sycl_kernel_entry_point`` attribute facilitates the generation of an
+offload kernel entry point, sometimes called a SYCL kernel caller function,
+suitable for invoking a SYCL kernel on an offload device. The attribute is
+intended for use in the implementation of SYCL kernel invocation functions
+like the ``single_task`` and ``parallel_for`` member functions of the
+``sycl::handler`` class specified in section 4.9.4, "Command group ``handler``
+class", of the SYCL 2020 specification.
+
+The attribute requires a single type argument that specifies a class type that
+meets the requirements for a SYCL kernel name as described in section 5.2,
+"Naming of kernels", of the SYCL 2020 specification. A unique kernel name type
+is required for each function declared with the attribute. The attribute may
+not first appear on a declaration that follows a definition of the function.
+
+The attribute only appertains to functions and only those that meet the
+following requirements.
+
+* Has a ``void`` return type.
+* Is not a non-static member function, constructor, or destructor.
+* Is not a C variadic function.
+* Is not a coroutine.
+* Is not defined as deleted or as defaulted.
+* Is not declared with the ``constexpr`` or ``consteval`` specifiers.
+* Is not declared with the ``[[noreturn]]`` attribute.
+
+Use in the implementation of a SYCL kernel invocation function might look as
+follows.
+
+.. code-block:: c++
+
+ namespace sycl {
+ class handler {
+ template<typename KernelNameType, typename KernelType>
+ [[ clang::sycl_kernel_entry_point(KernelNameType) ]]
+ static void kernel_entry_point(KernelType kernel) {
+ kernel();
+ }
+
+ public:
+ template<typename KernelNameType, typename KernelType>
+ void single_task(KernelType kernel) {
+ // Call kernel_entry_point() to trigger generation of an offload
+ // kernel entry point.
+ kernel_entry_point<KernelNameType>(kernel);
+ // Call functions appropriate for the desired offload backend
+ // (OpenCL, CUDA, HIP, Level Zero, etc...).
+ }
+ };
+ } // namespace sycl
+
+A SYCL kernel is a callable object of class type that is constructed on a host,
+often via a lambda expression, and then passed to a SYCL kernel invocation
+function to be executed on an offload device. A SYCL kernel invocation function
+is responsible for copying the provided SYCL kernel object to an offload
+device and initiating a call to it. The SYCL kernel object and its data members
+constitute the parameters of an offload kernel.
+
+A SYCL kernel type is required to satisfy the device copyability requirements
+specified in section 3.13.1, "Device copyable", of the SYCL 2020 specification.
+Additionally, any data members of the kernel object type are required to satisfy
+section 4.12.4, "Rules for parameter passing to kernels". For most types, these
+rules require that the type is trivially copyable. However, the SYCL
+specification mandates that certain special SYCL types, such as
+``sycl::accessor`` and ``sycl::stream`` be device copyable even if they are not
+trivially copyable. These types require special handling because they cannot
+be copied to device memory as if by ``memcpy()``. Additionally, some offload
+backends, OpenCL for example, require objects of some of these types to be
+passed as individual arguments to the offload kernel.
+
+An offload kernel consists of an entry point function that declares the
+parameters of the offload kernel and the set of all functions and variables that
+are directly or indirectly used by the entry point function.
+
+A SYCL kernel invocation function invokes a SYCL kernel on a device by
+performing the following tasks (likely with the help of an offload backend
+like OpenCL):
+
+#. Identifying the offload kernel entry point to be used for the SYCL kernel.
+
+#. Deconstructing the SYCL kernel object, if necessary, to produce the set of
+ offload kernel arguments required by the offload kernel entry point.
+
+#. Copying the offload kernel arguments to device memory.
+
+#. Initiating execution of the offload kernel entry point.
+
+The offload kernel entry point for a SYCL kernel performs the following tasks:
+
+#. Reconstituting the SYCL kernel object, if necessary, using the offload
+ kernel parameters.
+
+#. Calling the ``operator()`` member function of the (reconstituted) SYCL kernel
+ object.
+
+The ``sycl_kernel_entry_point`` attribute automates generation of an offload
+kernel entry point that performs those latter tasks. The parameters and body of
+a function declared with the ``sycl_kernel_entry_point`` attribute specify a
+pattern from which the parameters and body of the entry point function are
+derived. Consider the following call to a SYCL kernel invocation function.
+
+.. code-block:: c++
+
+ struct S { int i; };
+ void f(sycl::handler &handler, sycl::stream &sout, S s) {
+ handler.single_task<struct KN>([=] {
+ sout << "The value of s.i is " << s.i << "\n";
+ });
+ }
+
+The SYCL kernel object is the result of the lambda expression. It has two
+data members corresponding to the captures of ``sout`` and ``s``. Since one
+of these data members corresponds to a special SYCL type that must be passed
+individually as an offload kernel parameter, it is necessary to decompose the
+SYCL kernel object into its constituent parts; the offload kernel will have
+two kernel parameters. Given a SYCL implementation that uses a
+``sycl_kernel_entry_point`` attributed function like the one shown above, an
+offload kernel entry point function will be generated that looks approximately
+as follows.
+
+.. code-block:: c++
+
+ void sycl-kernel-caller-for-KN(sycl::stream sout, S s) {
+ kernel-type kernel = { sout, s );
+ kernel();
+ }
+
+There are a few items worthy of note:
+
+#. The name of the generated function incorporates the SYCL kernel name,
+ ``KN``, that was passed as the ``KernelNameType`` template parameter to
+ ``kernel_entry_point()`` and provided as the argument to the
+ ``sycl_kernel_entry_point`` attribute. There is a one-to-one correspondence
+ between SYCL kernel names and offload kernel entry points.
+
+#. The SYCL kernel is a lambda closure type and therefore has no name;
+ ``kernel-type`` is substituted above and corresponds to the ``KernelType``
+ template parameter deduced in the call to ``kernel_entry_point()``.
+ Lambda types cannot be declared and initialized using the aggregate
+ initialization syntax used above, but the intended behavior should be clear.
+
+#. ``S`` is a device copyable type that does not directly or indirectly contain
+ a data member of a SYCL special type. It therefore does not need to be
+ decomposed into its constituent members to be passed as a kernel argument.
+
+#. The depiction of the ``sycl::stream`` parameter as a single self contained
+ kernel parameter is an oversimplification. SYCL special types may require
+ additional decomposition such that the generated function might have three
+ or more parameters depending on how the SYCL library implementation defines
+ these types.
+
+#. The call to ``kernel_entry_point()`` has no effect other than to trigger
+ emission of the entry point function. The statments that make up the body
+ of the function are not executed when the function is called; they are
+ only used in the generation of the entry point function.
+
+It is not necessary for a function declared with the ``sycl_kernel_entry_point``
+attribute to be called for the offload kernel entry point to be emitted. For
+inline functions and function templates, any ODR-use will suffice. For other
+functions, an ODR-use is not required; the offload kernel entry point will be
+emitted if the function is defined.
+
+Functions declared with the ``sycl_kernel_entry_point`` attribute are not
+limited to the simple example shown above. They may have additional template
+parameters, declare additional function parameters, and have complex control
+flow in the function body. Function parameter decomposition and reconstitution
+is performed for all function parameters. The function must abide by the
+language feature restrictions described in section 5.4, "Language restrictions
+for device functions" in the SYCL 2020 specification.
+ }];
+}
+
def SYCLSpecialClassDocs : Documentation {
let Category = DocCatStmt;
let Content = [{
diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h
index 27c42b54018307..c9f3358124eda7 100644
--- a/clang/include/clang/Sema/SemaSYCL.h
+++ b/clang/include/clang/Sema/SemaSYCL.h
@@ -62,6 +62,7 @@ class SemaSYCL : public SemaBase {
ParsedType ParsedTy);
void handleKernelAttr(Decl *D, const ParsedAttr &AL);
+ void handleKernelEntryPointAttr(Decl *D, const ParsedAttr &AL);
};
} // namespace clang
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index 11e79d296cbec3..061a6866fb5fb1 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -14411,6 +14411,32 @@ void ASTContext::getFunctionFeatureMap(llvm::StringMap<bool> &FeatureMap,
}
}
+static SYCLKernelInfo BuildSYCLKernelInfo(CanQualType KernelNameType,
+ const FunctionDecl *FD) {
+ return {KernelNameType, FD};
+}
+
+void ASTContext::registerSYCLEntryPointFunction(FunctionDecl *FD) {
+ // If the function declaration to register is invalid or dependent, the
+ // registration attempt is ignored.
+ if (FD->isInvalidDecl() || FD->isTemplated())
+ return;
+
+ const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>();
+ assert(SKEPAttr && "Missing sycl_kernel_entry_point attribute");
+
+ // Be tolerant of multiple registration attempts so long as each attempt
+ // is for the same entity. Callers are obligated to detect and diagnose
+ // conflicting kernel names prior to calling this function.
+ CanQualType KernelNameType = getCanonicalType(SKEPAttr->getKernelName());
+ auto IT = SYCLKernels.find(KernelNameType);
+ assert((IT == SYCLKernels.end() ||
+ declaresSameEntity(FD, IT->second.getKernelEntryPointDecl())) &&
+ "SYCL kernel name conflict");
+ SYCLKernels.insert(
+ std::make_pair(KernelNameType, BuildSYCLKernelInfo(KernelNameType, FD)));
+}
+
OMPTraitInfo &ASTContext::getNewOMPTraitInfo() {
OMPTraitInfoVector.emplace_back(new OMPTraitInfo());
return *OMPTraitInfoVector.back();
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 1aa3e8edfe1b13..00c8f871bbb38c 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -12094,6 +12094,9 @@ bool Sema::CheckFunctionDeclaration(Scope *S, FunctionDecl *NewFD,
if (LangOpts.OpenMP)
OpenMP().ActOnFinishedFunctionDefinitionInOpenMPAssumeScope(NewFD);
+ if (LangOpts.isSYCL() && NewFD->hasAttr<SYCLKernelEntryPointAttr>())
+ getASTContext().registerSYCLEntryPointFunction(NewFD);
+
// Semantic checking for this function declaration (in isolation).
if (getLangOpts().CPlusPlus) {
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 601c6f2eef1d9c..a90f8706fc0080 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -6620,6 +6620,9 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL,
case ParsedAttr::AT_SYCLKernel:
S.SYCL().handleKernelAttr(D, AL);
break;
+ case ParsedAttr::AT_SYCLKernelEntryPoint:
+ S.SYCL().handleKernelEntryPointAttr(D, AL);
+ break;
case ParsedAttr::AT_SYCLSpecialClass:
handleSimpleAttribute<SYCLSpecialClassAttr>(S, D, AL);
break;
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index f2d13d456c25fc..e7cecebae25808 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -198,3 +198,12 @@ void SemaSYCL::handleKernelAttr(Decl *D, const ParsedAttr &AL) {
handleSimpleAttribute<SYCLKernelAttr>(*this, D, AL);
}
+
+void SemaSYCL::handleKernelEntryPointAttr(Decl *D, const ParsedAttr &AL) {
+ ParsedType PT = AL.getTypeArg();
+ TypeSourceInfo *TSI = nullptr;
+ (void)SemaRef.GetTypeFromParser(PT, &TSI);
+ assert(TSI && "no type source info for attribute argument");
+ D->addAttr(::new (SemaRef.Context)
+ SYCLKernelEntryPointAttr(SemaRef.Context, AL, TSI));
+}
diff --git a/clang/lib/Serialization/ASTReaderDecl.cpp b/clang/lib/Serialization/ASTReaderDecl.cpp
index d4e392dcc6bcd0..20edd53598e5bd 100644
--- a/clang/lib/Serialization/ASTReaderDecl.cpp
+++ b/clang/lib/Serialization/ASTReaderDecl.cpp
@@ -1155,6 +1155,14 @@ void ASTDeclReader::VisitFunctionDecl(FunctionDecl *FD) {
for (unsigned I = 0; I != NumParams; ++I)
Params.push_back(readDeclAs<ParmVarDecl>());
FD->setParams(Reader.getContext(), Params);
+
+ // If the declaration is a SYCL kernel entry point function as indicated by
+ // the presence of a sycl_kernel_entry_point attribute, register it so that
+ // associated metadata is recreated.
+ if (FD->hasAttr<SYCLKernelEntryPointAttr>()) {
+ ASTContext &C = Reader.getContext();
+ C.registerSYCLEntryPointFunction(FD);
+ }
}
void ASTDeclReader::VisitObjCMethodDecl(ObjCMethodDecl *MD) {
diff --git a/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp b/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
new file mode 100644
index 00000000000000..c351f3b7d03eab
--- /dev/null
+++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
@@ -0,0 +1,144 @@
+// Tests without serialization:
+// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-unknown -fsycl-is-device \
+// RUN: -ast-dump %s \
+// RUN: | FileCheck --match-full-lines %s
+// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-unknown -fsycl-is-host \
+// RUN: -ast-dump %s \
+// RUN: | FileCheck --match-full-lines %s
+//
+// Tests with serialization:
+// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-unknown -fsycl-is-device \
+// RUN: -emit-pch -o %t %s
+// RUN: %clang_cc1 -x c++ -std=c++17 -triple x86_64-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 --match-full-lines %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 --match-full-lines %s
+
+// These tests validate the AST produced for functions declared with the
+// sycl_kernel_entry_point attribute.
+
+// CHECK: TranslationUnitDecl {{.*}}
+
+// A unique kernel name type is required for each declared kernel entry point.
+template<int, int=0> struct KN;
+
+__attribute__((sycl_kernel_entry_point(KN<1>)))
+void skep1() {
+}
+// CHECK: |-FunctionDecl {{.*}} skep1 'void ()'
+// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} KN<1>
+
+using KN2 = KN<2>;
+__attribute__((sycl_kernel_entry_point(KN2)))
+void skep2() {
+}
+// CHECK: |-FunctionDecl {{.*}} skep2 'void ()'
+// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} KN2
+
+template<int I> using KNT = KN<I>;
+__attribute__((sycl_kernel_entry_point(KNT<3>)))
+void skep3() {
+}
+// CHECK: |-FunctionDecl {{.*}} skep3 'void ()'
+// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} KNT<3>
+
+template<typename KNT, typename F>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void skep4(F f) {
+ f();
+}
+// CHECK: |-FunctionTemplateDecl {{.*}} skep4
+// CHECK-NEXT: | |-TemplateTypeParmDecl {{.*}} KNT
+// CHECK-NEXT: | |-TemplateTypeParmDecl {{.*}} F
+// CHECK-NEXT: | |-FunctionDecl {{.*}} skep4 'void (F)'
+// CHECK: | | `-SYCLKernelEntryPointAttr {{.*}} KNT
+
+void test_skep4() {
+ skep4<KNT<4>>([]{});
+}
+// CHECK: | `-FunctionDecl {{.*}} used skep4 'void ((lambda at {{.*}}))' implicit_instantiation
+// CHECK-NEXT: | |-TemplateArgument type 'KN<4>'
+// CHECK: | |-TemplateArgument type '(lambda at {{.*}})'
+// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} struct KN<4>
+// CHECK-NEXT: |-FunctionDecl {{.*}} test_skep4 'void ()'
+
+template<typename KNT, typename T>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void skep5(T) {
+}
+// CHECK: |-FunctionTemplateDecl {{.*}} skep5
+// CHECK-NEXT: | |-TemplateTypeParmDecl {{.*}} KNT
+// CHECK-NEXT: | |-TemplateTypeParmDecl {{.*}} T
+// CHECK-NEXT: | |-FunctionDecl {{.*}} skep5 'void (T)'
+// CHECK: | | `-SYCLKernelEntryPointAttr {{.*}} KNT
+
+// Checks for the explicit template instantiation declaration below.
+// CHECK: | `-FunctionDecl {{.*}} skep5 'void (int)' explicit_instantiation_definition
+// CHECK-NEXT: | |-TemplateArgument type 'KN<5, 4>'
+// CHECK: | |-TemplateArgument type 'int'
+// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} KN<5, 4>
+
+// 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.
+// FIXME: Clang currently instantiates and propagates attributes from a function
+// FIXME: template to its explicit specializations resulting in the following
+// FIXME: explicit specialization having an attribute incorrectly attached.
+template<>
+void skep5<KN<5,1>>(short) {
+}
+// CHECK: |-FunctionDecl {{.*}} prev {{.*}} skep5 'void (short)' explicit_specialization
+// CHECK-NEXT: | |-TemplateArgument type 'KN<5, 1>'
+// CHECK: | |-TemplateArgument type 'short'
+// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} Inherited struct KN<5, 1>
+
+template<>
+[[clang::sycl_kernel_entry_point(KN<5,2>)]]
+void skep5<KN<5,2>>(long) {
+}
+// CHECK: |-FunctionDecl {{.*}} prev {{.*}} skep5 'void (long)' explicit_specialization
+// CHECK-NEXT: | |-TemplateArgument type 'KN<5, 2>'
+// CHECK: | |-TemplateArgument type 'long'
+// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} KN<5, 2>
+
+template<>
+[[clang::sycl_kernel_entry_point(KN<5,3>)]]
+void skep5<KN<5,-1>>(long long) {
+}
+// CHECK: |-FunctionDecl {{.*}} prev {{.*}} skep5 'void (long long)' explicit_specialization
+// CHECK-NEXT: | |-TemplateArgument type 'KN<5, -1>'
+// CHECK: | |-TemplateArgument type 'long long'
+// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} KN<5, 3>
+
+template void skep5<KN<5,4>>(int);
+// Checks are located with the primary template declaration above.
+
+// Ensure that matching attributes from multiple declarations are ok.
+[[clang::sycl_kernel_entry_point(KN<6>)]]
+void skep6();
+[[clang::sycl_kernel_entry_point(KN<6>)]]
+void skep6() {
+}
+// CHECK: |-FunctionDecl {{.*}} skep6 'void ()'
+// CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} KN<6>
+// CHECK-NEXT: |-FunctionDecl {{.*}} prev {{.*}} skep6 'void ()'
+// CHECK-NEXT: | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} KN<6>
+
+// Ensure that matching attributes from the same declaration are ok.
+[[clang::sycl_kernel_entry_point(KN<7>), clang::sycl_kernel_entry_point(KN<7>)]]
+void skep7() {
+}
+// CHECK: |-FunctionDecl {{.*}} skep7 'void ()'
+// CHECK-NEXT: | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | |-SYCLKernelEntryPointAttr {{.*}} KN<7>
+// CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} KN<7>
+
+void the_end() {}
+// CHECK: `-FunctionDecl {{.*}} the_end 'void ()'
diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
index e28b0775410c0a..3f1977d4d408b9 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -178,6 +178,7 @@
// CHECK-NEXT: ReturnTypestate (SubjectMatchRule_function, SubjectMatchRule_variable_is_parameter)
// CHECK-NEXT: ReturnsNonNull (SubjectMatchRule_objc_method, SubjectMatchRule_function)
// CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function)
+// CHECK-NEXT: SYCLKernelEntryPoint (SubjectMatchRule_function)
// CHECK-NEXT: SYCLSpecialClass (SubjectMatchRule_record)
// CHECK-NEXT: ScopedLockable (SubjectMatchRule_record)
// CHECK-NEXT: Section (SubjectMatchRule_function, SubjectMatchRule_variable_is_global, SubjectMatchRule_objc_method, SubjectMatchRule_objc_property)
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
new file mode 100644
index 00000000000000..c63d241163e618
--- /dev/null
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
@@ -0,0 +1,137 @@
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-device -verify %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only -fsycl-is-device -verify %s
+
+// These tests validate parsing of the sycl_kernel_entry_point argument list
+// and that the single argument names a type.
+
+// Templates used to exercise class template specializations.
+template<int> struct ST; // #ST-decl
+template<int N> using TTA = ST<N>; // #TTA-decl
+
+
+////////////////////////////////////////////////////////////////////////////////
+// Valid declarations.
+////////////////////////////////////////////////////////////////////////////////
+
+struct S1;
+[[clang::sycl_kernel_entry_point(S1)]] void ok1();
+
+typedef struct {} TA2;
+[[clang::sycl_kernel_entry_point(TA2)]] void ok2();
+
+using TA3 = struct {};
+[[clang::sycl_kernel_entry_point(TA3)]] void ok3();
+
+[[clang::sycl_kernel_entry_point(ST<4>)]] void ok4();
+
+[[clang::sycl_kernel_entry_point(TTA<5>)]] void ok5();
+
+namespace NS6 {
+ struct NSS;
+}
+[[clang::sycl_kernel_entry_point(NS6::NSS)]] void ok6();
+
+namespace {
+ struct UNSS7;
+}
+[[clang::sycl_kernel_entry_point(UNSS7)]] void ok7();
+
+struct {} s;
+[[clang::sycl_kernel_entry_point(decltype(s))]] void ok8();
+
+template<typename KN>
+[[clang::sycl_kernel_entry_point(KN)]] void ok9();
+void test_ok9() {
+ ok9<struct LS1>();
+}
+
+template<int, typename KN>
+[[clang::sycl_kernel_entry_point(KN)]] void ok10();
+void test_ok10() {
+ ok10<1, struct LS2>();
+}
+
+namespace NS11 {
+ struct NSS;
+}
+template<typename T>
+[[clang::sycl_kernel_entry_point(T)]] void ok11() {}
+template<>
+[[clang::sycl_kernel_entry_point(NS11::NSS)]] void ok11<NS11::NSS>() {}
+
+struct S12;
+[[clang::sycl_kernel_entry_point(S12)]] void ok12();
+[[clang::sycl_kernel_entry_point(S12)]] void ok12() {}
+
+template<typename T>
+[[clang::sycl_kernel_entry_point(T)]] void ok13(T k);
+void test_ok13() {
+ ok13([]{});
+}
+
+
+////////////////////////////////////////////////////////////////////////////////
+// Invalid declarations.
+////////////////////////////////////////////////////////////////////////////////
+
+// expected-error at +1 {{'sycl_kernel_entry_point' attribute takes one argument}}
+[[clang::sycl_kernel_entry_point]] void bad1();
+
+// expected-error at +1 {{'sycl_kernel_entry_point' attribute takes one argument}}
+[[clang::sycl_kernel_entry_point()]] void bad2();
+
+struct B3;
+// expected-error at +2 {{expected ')'}}
+// expected-error at +1 {{expected ']'}}
+[[clang::sycl_kernel_entry_point(B3,)]] void bad3();
+
+struct B4;
+// expected-error at +3 {{expected ')'}}
+// expected-error at +2 {{expected ','}}
+// expected-warning at +1 {{unknown attribute 'X' ignored}}
+[[clang::sycl_kernel_entry_point(B4, X)]] void bad4();
+
+// expected-error at +1 {{expected a type}}
+[[clang::sycl_kernel_entry_point(1)]] void bad5();
+
+void f6();
+// expected-error at +1 {{unknown type name 'f6'}}
+[[clang::sycl_kernel_entry_point(f6)]] void bad6();
+
+// expected-error at +2 {{use of class template 'ST' requires template arguments; argument deduction not allowed here}}
+// expected-note@#ST-decl {{template is declared here}}
+[[clang::sycl_kernel_entry_point(ST)]] void bad7();
+
+// expected-error at +2 {{use of alias template 'TTA' requires template arguments; argument deduction not allowed here}}
+// expected-note@#TTA-decl {{template is declared here}}
+[[clang::sycl_kernel_entry_point(TTA)]] void bad8();
+
+enum {
+ e9
+};
+// expected-error at +1 {{unknown type name 'e9'}}
+[[clang::sycl_kernel_entry_point(e9)]] void bad9();
+
+#if __cplusplus >= 202002L
+template<typename> concept C = true;
+// expected-error at +1 {{expected a type}}
+[[clang::sycl_kernel_entry_point(C)]] void bad10();
+
+// expected-error at +1 {{expected a type}}
+[[clang::sycl_kernel_entry_point(C<int>)]] void bad11();
+#endif
+
+struct B12; // #B12-decl
+// 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.
+// FIXME: Clang currently instantiates and propagates attributes from a function
+// FIXME: template to its explicit specializations resulting in the following
+// FIXME: spurious error.
+// expected-error at +4 {{incomplete type 'B12' named in nested name specifier}}
+// expected-note at +5 {{in instantiation of function template specialization 'bad12<B12>' requested here}}
+// expected-note@#B12-decl {{forward declaration of 'B12'}}
+template<typename T>
+[[clang::sycl_kernel_entry_point(typename T::not_found)]] void bad12() {}
+template<>
+void bad12<B12>() {}
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-ignored.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-ignored.cpp
new file mode 100644
index 00000000000000..30de6ae0b0e6f9
--- /dev/null
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-ignored.cpp
@@ -0,0 +1,17 @@
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -verify %s
+
+// These tests validate that the sycl_kernel_entry_point attribute is ignored
+// when SYCL support is not enabled.
+
+// A unique kernel name type is required for each declared kernel entry point.
+template<int> struct KN;
+
+// expected-warning at +1 {{'sycl_kernel_entry_point' attribute ignored}}
+[[clang::sycl_kernel_entry_point(KN<1>)]]
+void ok1();
+
+// expected-warning at +2 {{'sycl_kernel_entry_point' attribute ignored}}
+template<typename KNT>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void ok2() {}
+template void ok2<KN<2>>();
More information about the cfe-commits
mailing list