diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h index b65a1f7dff5bc1..eae80aed479f8d 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 { /// in device compilation. llvm::DenseSet 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 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,14 @@ class ASTContext : public RefCountedBase { void getFunctionFeatureMap(llvm::StringMap &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 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..87782146ae14ba 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 + [[ clang::sycl_kernel_entry_point(KernelNameType) ]] + static void kernel_entry_point(KernelType kernel) { + kernel(); + } + + public: + template + void single_task(KernelType kernel) { + // Call kernel_entry_point() to trigger generation of an offload + // kernel entry point. + kernel_entry_point(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([=] { + 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 85b3984940ffc2..846a1ab45a9fd0 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -14296,6 +14296,32 @@ void ASTContext::getFunctionFeatureMap(llvm::StringMap &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(); + 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 de8805e15bc750..2ce868c403c693 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -12053,6 +12053,9 @@ bool Sema::CheckFunctionDeclaration(Scope *S, FunctionDecl *NewFD, if (LangOpts.OpenMP) OpenMP().ActOnFinishedFunctionDefinitionInOpenMPAssumeScope(NewFD); + if (LangOpts.isSYCL() && NewFD->hasAttr()) + 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(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(*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..c8a59e9b929917 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()); 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()) { + 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/ //" -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/ //" -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 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 using KNT = KN; +__attribute__((sycl_kernel_entry_point(KNT<3>))) +void skep3() { +} +// CHECK: |-FunctionDecl {{.*}} skep3 'void ()' +// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} KNT<3> + +template +[[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>([]{}); +} +// 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 +[[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>(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>(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>(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>(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 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 struct ST; // #ST-decl +template using TTA = ST; // #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 +[[clang::sycl_kernel_entry_point(KN)]] void ok9(); +void test_ok9() { + ok9(); +} + +template +[[clang::sycl_kernel_entry_point(KN)]] void ok10(); +void test_ok10() { + ok10<1, struct LS2>(); +} + +namespace NS11 { + struct NSS; +} +template +[[clang::sycl_kernel_entry_point(T)]] void ok11() {} +template<> +[[clang::sycl_kernel_entry_point(NS11::NSS)]] void ok11() {} + +struct S12; +[[clang::sycl_kernel_entry_point(S12)]] void ok12(); +[[clang::sycl_kernel_entry_point(S12)]] void ok12() {} + +template +[[clang::sycl_kernel_entry_point(T)]] void ok13(T k); +void test_ok13() { + ok13([]{}); +} + + +//////////////////////////////////////////////////////////////////////////////// +// Invalid declarations. +//////////////////////////////////////////////////////////////////////////////// + +// expected-error@+1 {{'sycl_kernel_entry_point' attribute takes one argument}} +[[clang::sycl_kernel_entry_point]] void bad1(); + +// expected-error@+1 {{'sycl_kernel_entry_point' attribute takes one argument}} +[[clang::sycl_kernel_entry_point()]] void bad2(); + +struct B3; +// expected-error@+2 {{expected ')'}} +// expected-error@+1 {{expected ']'}} +[[clang::sycl_kernel_entry_point(B3,)]] void bad3(); + +struct B4; +// expected-error@+3 {{expected ')'}} +// expected-error@+2 {{expected ','}} +// expected-warning@+1 {{unknown attribute 'X' ignored}} +[[clang::sycl_kernel_entry_point(B4, X)]] void bad4(); + +// expected-error@+1 {{expected a type}} +[[clang::sycl_kernel_entry_point(1)]] void bad5(); + +void f6(); +// expected-error@+1 {{unknown type name 'f6'}} +[[clang::sycl_kernel_entry_point(f6)]] void bad6(); + +// expected-error@+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@+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@+1 {{unknown type name 'e9'}} +[[clang::sycl_kernel_entry_point(e9)]] void bad9(); + +#if __cplusplus >= 202002L +template concept C = true; +// expected-error@+1 {{expected a type}} +[[clang::sycl_kernel_entry_point(C)]] void bad10(); + +// expected-error@+1 {{expected a type}} +[[clang::sycl_kernel_entry_point(C)]] 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@+4 {{incomplete type 'B12' named in nested name specifier}} +// expected-note@+5 {{in instantiation of function template specialization 'bad12' requested here}} +// expected-note@#B12-decl {{forward declaration of 'B12'}} +template +[[clang::sycl_kernel_entry_point(typename T::not_found)]] void bad12() {} +template<> +void bad12() {} 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 struct KN; + +// expected-warning@+1 {{'sycl_kernel_entry_point' attribute ignored}} +[[clang::sycl_kernel_entry_point(KN<1>)]] +void ok1(); + +// expected-warning@+2 {{'sycl_kernel_entry_point' attribute ignored}} +template +[[clang::sycl_kernel_entry_point(KNT)]] +void ok2() {} +template void ok2>();