[SYCL] The sycl_kernel_entry_point attribute. #111389
Merged
+591
−3
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
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.
tahonermann
requested review from
bader,
AaronBallman,
yxsamliu,
jdoerfert,
erichkeane,
Endilll and
elizabethandrews
llvmbot
added
clang
|
@llvm/pr-subscribers-clang Author: Tom Honermann (tahonermann) ChangesThe Properties of the offload kernel are collected when a function declared with the The 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. Patch is 24.90 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/111389.diff 14 Files Affected:
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
+
+
+//////////////////////////////////////////////////////////...
[truncated]
|
|
@llvm/pr-subscribers-clang-modules Author: Tom Honermann (tahonermann) ChangesThe Properties of the offload kernel are collected when a function declared with the The 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. Patch is 24.90 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/111389.diff 14 Files Affected:
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
+
+
+//////////////////////////////////////////////////////////...
[truncated]
|
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
|
This is the first in a series of PRs to provide core language support for SYCL-aware host and device compilation and support for SYCL library implementations. The following links present a preliminary set of changes for future PRs that depend on this one. Some of these need a little work yet (e.g., some of the diagnostics need a bit of tweaking), but will be submitted in order as dependencies are accepted.
The last one listed above is to remove the existing |
erichkeane
reviewed
Oct 7, 2024
|
@erichkeane, I think I have addressed or replied to all of your comments. Another round of review feedback would be much appreciated! @bader, as SYCL code maintainer, I would appreciate if you can take a look. |
I didnt get a chance to get back to this during the week, but intend to after llvm-dev. By the time you got to responding, I'd context-switched it, so I'll have to spend some time getting back to it. |
keryell
approved these changes
Oct 21, 2024
erichkeane
reviewed
Oct 28, 2024
erichkeane
requested changes
Oct 30, 2024
Changes include: - Removal of an unneeded inclusion of <string>. - An update to the `sycl_kernel_entry_point` attribute documentation. - Updates to the AST test to validate next line proximity where applicable.
Changes include: - Use `isTemplated()` instead of `isDependentContext()` since the latter doesn't necessarily handle templated friend functions.
Changes include: - Replaced a use of `llvm::report_fatal_error()` with `assert()`. - Modified `registerSYCLEntryPointFunction()` to silently skip registration of invalid or dependent function declarations.
|
@erichkeane, I believe all of the concerns you raised have been resolved with the exception of the |
erichkeane
approved these changes
Nov 4, 2024
There was a problem hiding this comment.
1 nit, else LGTM.
Please let @keryell chime in on the constexpr decision.
Changes include: - Minor refactoring to eliminate a branch that only contained an assertion.
PhilippRados
pushed a commit
to PhilippRados/llvm-project
that referenced
this pull request
Nov 6, 2024
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.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
The
sycl_kernel_entry_pointattribute 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_pointattribute 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_pointattribute is intended to replace the existingsycl_kernelattribute 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.