[SYCL] The sycl_kernel_entry_point attribute.

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
   //===--------------------------------------------------------------------===//

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

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];
 }
 

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 = [{

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

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();

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) {

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;

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));
+}

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) {