[clang] [SYCL] The sycl_kernel_entry_point attribute. (PR #111389)
Tom Honermann via cfe-commits
cfe-commits at lists.llvm.org
Mon Oct 7 07:57:11 PDT 2024
https://github.com/tahonermann created https://github.com/llvm/llvm-project/pull/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](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:naming.kernels).
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.
>From 3c4a2b8a52d3f1c730df88a308dece21a67834ef Mon Sep 17 00:00:00 2001
From: Tom Honermann <tom.honermann at intel.com>
Date: Fri, 4 Oct 2024 11:10:32 -0700
Subject: [PATCH] [SYCL] The sycl_kernel_entry_point attribute.
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.
---
clang/include/clang/AST/ASTContext.h | 12 ++
clang/include/clang/AST/SYCLKernelInfo.h | 47 ++++++
clang/include/clang/Basic/Attr.td | 16 +-
clang/include/clang/Basic/AttrDocs.td | 58 +++++++
clang/include/clang/Sema/SemaSYCL.h | 1 +
clang/lib/AST/ASTContext.cpp | 25 +++
clang/lib/Sema/SemaDecl.cpp | 4 +
clang/lib/Sema/SemaDeclAttr.cpp | 3 +
clang/lib/Sema/SemaSYCL.cpp | 9 ++
clang/lib/Serialization/ASTReaderDecl.cpp | 9 ++
.../ast-dump-sycl-kernel-entry-point.cpp | 144 ++++++++++++++++++
...a-attribute-supported-attributes-list.test | 1 +
.../sycl-kernel-entry-point-attr-grammar.cpp | 137 +++++++++++++++++
.../sycl-kernel-entry-point-attr-ignored.cpp | 17 +++
14 files changed, 480 insertions(+), 3 deletions(-)
create mode 100644 clang/include/clang/AST/SYCLKernelInfo.h
create mode 100644 clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
create mode 100644 clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
create mode 100644 clang/test/SemaSYCL/sycl-kernel-entry-point-attr-ignored.cpp
diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h
index b65a1f7dff5bc1..e68f0a4da57b16 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"
@@ -1222,6 +1223,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
@@ -3301,6 +3307,12 @@ 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.
+ 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..79a83330f1d23b
--- /dev/null
+++ b/clang/include/clang/AST/SYCLKernelInfo.h
@@ -0,0 +1,47 @@
+//===--- 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 <string>
+#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 ce86116680d7a3..c4a3615752bf10 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">;
@@ -1489,14 +1490,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 8ef151b3f2fddb..cd3aec8f70f024 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -455,6 +455,64 @@ 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 specifies that a function definition
+defines a pattern for an offload kernel entry point function to be emitted when
+the source code is compiled with ``-fsycl`` for a device target. Such functions
+serve as the execution entry point for a SYCL run-time library to invoke a SYCL
+kernel on a device. The function's parameters define the parameters to the
+offload kernel.
+
+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 appertains only to non-member functions and static member
+functions that meet the following requirements:
+
+- Has a ``void`` return type.
+- Is not a 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.
+
+This attribute is intended for use in the implementation of SYCL run-time
+libraries that implement 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. Such use might look something like the following.
+
+.. 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) {
+ kernel_entry_point<KernelNameType>(kernel);
+ }
+ };
+ } // namespace sycl
+
+It is not necessary for a SYCL kernel entry point function 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.
+ }];
+}
+
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 85b3984940ffc2..daf4b8398bdd55 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -14296,6 +14296,31 @@ void ASTContext::getFunctionFeatureMap(llvm::StringMap<bool> &FeatureMap,
}
}
+static SYCLKernelInfo BuildSYCLKernelInfo(ASTContext &Context,
+ CanQualType KernelNameType,
+ const FunctionDecl *FD) {
+ return { KernelNameType, FD };
+}
+
+void ASTContext::registerSYCLEntryPointFunction(FunctionDecl *FD) {
+ assert(!FD->isInvalidDecl());
+ assert(!FD->isDependentContext());
+
+ const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>();
+ assert(SKEPAttr && "Missing sycl_kernel_entry_point attribute");
+
+ CanQualType KernelNameType = getCanonicalType(SKEPAttr->getKernelName());
+ auto IT = SYCLKernels.find(KernelNameType);
+ if (IT != SYCLKernels.end()) {
+ if (!declaresSameEntity(FD, IT->second.GetKernelEntryPointDecl()))
+ llvm::report_fatal_error("SYCL kernel name conflict");
+ } else {
+ SYCLKernels.insert_or_assign(
+ KernelNameType,
+ BuildSYCLKernelInfo(*this, 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 de8805e15bc750..6af5264f79e1f9 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -12053,6 +12053,10 @@ bool Sema::CheckFunctionDeclaration(Scope *S, FunctionDecl *NewFD,
if (LangOpts.OpenMP)
OpenMP().ActOnFinishedFunctionDefinitionInOpenMPAssumeScope(NewFD);
+ if (LangOpts.isSYCL() && NewFD->hasAttr<SYCLKernelEntryPointAttr>() &&
+ !NewFD->isInvalidDecl() && !NewFD->isDependentContext())
+ 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 14cc51cf89665a..2504fc2e7b8ec9 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -6606,6 +6606,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..a168e654ad6876 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 321e65fd2b094f..02e723aef0ff8a 100644
--- a/clang/lib/Serialization/ASTReaderDecl.cpp
+++ b/clang/lib/Serialization/ASTReaderDecl.cpp
@@ -1155,6 +1155,15 @@ 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->isInvalidDecl() && !FD->isDependentContext() &&
+ 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..342d2f71c357e2
--- /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: | |-TemplateTypeParmDecl {{.*}} KNT
+// CHECK: | |-TemplateTypeParmDecl {{.*}} F
+// CHECK: | |-FunctionDecl {{.*}} skep4 'void (F)'
+// CHECK: | | `-SYCLKernelEntryPointAttr {{.*}} KNT
+
+void test_skep4() {
+ skep4<KNT<4>>([]{});
+}
+// CHECK: | `-FunctionDecl {{.*}} used skep4 'void ((lambda at {{.*}}))' implicit_instantiation
+// CHECK: | |-TemplateArgument type 'KN<4>'
+// CHECK: | |-TemplateArgument type '(lambda at {{.*}})'
+// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} struct KN<4>
+// CHECK: |-FunctionDecl {{.*}} test_skep4 'void ()'
+
+template<typename KNT, typename T>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void skep5(T) {
+}
+// CHECK: |-FunctionTemplateDecl {{.*}} skep5
+// CHECK: | |-TemplateTypeParmDecl {{.*}} KNT
+// CHECK: | |-TemplateTypeParmDecl {{.*}} T
+// CHECK: | |-FunctionDecl {{.*}} skep5 'void (T)'
+// CHECK: | | `-SYCLKernelEntryPointAttr {{.*}} KNT
+
+// Checks for the explicit template instantiation declaration below.
+// CHECK: | `-FunctionDecl {{.*}} skep5 'void (int)' explicit_instantiation_definition
+// CHECK: | |-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: | |-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: | |-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: | |-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: |-FunctionDecl {{.*}} prev {{.*}} skep6 'void ()'
+// CHECK: | |-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: | |-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 914f94c08a9fd9..5c2f3a347dfb79 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -179,6 +179,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