Refactor diagnostics for sycl_kernel_entry_point

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,22 @@ 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);
+
+  /// Given a type used as a SYCL kernel name, returns a reference to the
+  /// metadata generated from the corresponding SYCL kernel entry point.
+  /// Aborts if the provided type is not a registered SYCL kernel name.
+  const SYCLKernelInfo &getSYCLKernelInfo(QualType T) const;
+
+  /// Returns a pointer to the metadata generated from the corresponding
+  /// SYCLkernel entry point if the provided type corresponds to a registered
+  /// SYCL kernel name. Returns a null pointer otherwise.
+  const SYCLKernelInfo *findSYCLKernelInfo(QualType T) const;
+
   //===--------------------------------------------------------------------===//
   //                    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/Basic/DiagnosticGroups.td

@@ -641,6 +641,7 @@ def PoundPragmaMessage : DiagGroup<"#pragma-messages">,
 def : DiagGroup<"redundant-decls">;
 def RedeclaredClassMember : DiagGroup<"redeclared-class-member">;
 def GNURedeclaredEnum : DiagGroup<"gnu-redeclared-enum">;
+def RedundantAttribute : DiagGroup<"redundant-attribute">;
 def RedundantMove : DiagGroup<"redundant-move">;
 def Register : DiagGroup<"register", [DeprecatedRegister]>;
 def ReturnTypeCLinkage : DiagGroup<"return-type-c-linkage">;

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12288,6 +12288,31 @@ def err_sycl_special_type_num_init_method : Error<
   "types with 'sycl_special_class' attribute must have one and only one '__init' "
   "method defined">;
 
+// SYCL kernel entry point diagnostics
+def err_sycl_entry_point_invalid : Error<
+  "'sycl_kernel_entry_point' attribute cannot be applied to a"
+  " %select{non-static member|variadic|deleted|defaulted|constexpr|consteval|"
+           "noreturn|coroutine}0 function">;
+def err_sycl_entry_point_invalid_redeclaration : Error<
+  "'sycl_kernel_entry_point' kernel name argument does not match prior"
+  " declaration%diff{: $ vs $|}0,1">;
+def err_sycl_kernel_name_conflict : Error<
+  "'sycl_kernel_entry_point' kernel name %0 conflicts with a previous"
+  " declaration">;
+def warn_sycl_kernel_name_not_a_class_type : Warning<
+  "%0 is not a valid SYCL kernel name type; a class type is required">,
+  InGroup<DiagGroup<"nonportable-sycl">>, DefaultError;
+def warn_sycl_entry_point_redundant_declaration : Warning<
+  "redundant 'sycl_kernel_entry_point' attribute">, InGroup<RedundantAttribute>;
+def err_sycl_entry_point_after_definition : Error<
+  "'sycl_kernel_entry_point' attribute cannot be added to a function after the"
+  " function is defined">;
+def err_sycl_entry_point_return_type : Error<
+  "'sycl_kernel_entry_point' attribute only applies to functions with a"
+  " 'void' return type">;
+def err_sycl_entry_point_on_main : Error<
+  "'main' cannot be declared with the 'sycl_kernel_entry_point' attribute">;
+
 def warn_cuda_maxclusterrank_sm_90 : Warning<
   "maxclusterrank requires sm_90 or higher, CUDA arch provided: %0, ignoring "
   "%1 attribute">, InGroup<IgnoredAttributes>;

clang/include/clang/Sema/Sema.h

@@ -13237,6 +13237,8 @@ class Sema final : public SemaBase {
   /// Prints the current instantiation stack through a series of
   /// notes.
   void PrintInstantiationStack();
+  void
+      PrintInstantiationStack(std::function<void(const PartialDiagnosticAt &)>);
 
   /// Determines whether we are currently in a context where
   /// template argument substitution failures are not considered

clang/include/clang/Sema/SemaSYCL.h

@@ -28,6 +28,12 @@ class SemaSYCL : public SemaBase {
 public:
   SemaSYCL(Sema &S);
 
+  using ContextNotes = SmallVector<PartialDiagnosticAt, 1>;
+  llvm::DenseMap<CanonicalDeclPtr<const FunctionDecl>, ContextNotes>
+      SYCLKernelEntryContextNotes;
+  llvm::DenseSet<CanonicalDeclPtr<const FunctionDecl>>
+      DiagnosedSYCLKernelEntryPoint;
+
   /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current
   /// context is "used as device code".
   ///
@@ -62,6 +68,9 @@ class SemaSYCL : public SemaBase {
                                        ParsedType ParsedTy);
 
   void handleKernelAttr(Decl *D, const ParsedAttr &AL);
+  void handleKernelEntryPointAttr(Decl *D, const ParsedAttr &AL);
+
+  void CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD, bool CheckUseOfDecl);
 };
 
 } // namespace clang

clang/lib/AST/ASTContext.cpp

@@ -14296,6 +14296,44 @@ 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));
+  }
+}
+
+const SYCLKernelInfo &ASTContext::getSYCLKernelInfo(QualType T) const {
+  CanQualType KernelNameType = getCanonicalType(T);
+  return SYCLKernels.at(KernelNameType);
+}
+
+const SYCLKernelInfo *ASTContext::findSYCLKernelInfo(QualType T) const {
+  CanQualType KernelNameType = getCanonicalType(T);
+  auto IT = SYCLKernels.find(KernelNameType);
+  if (IT != SYCLKernels.end())
+    return &IT->second;
+  return nullptr;
+}
+
 OMPTraitInfo &ASTContext::getNewOMPTraitInfo() {
   OMPTraitInfoVector.emplace_back(new OMPTraitInfo());
   return *OMPTraitInfoVector.back();

clang/lib/Sema/SemaDecl.cpp

@@ -54,6 +54,7 @@
 #include "clang/Sema/SemaPPC.h"
 #include "clang/Sema/SemaRISCV.h"
 #include "clang/Sema/SemaSwift.h"
+#include "clang/Sema/SemaSYCL.h"
 #include "clang/Sema/SemaWasm.h"
 #include "clang/Sema/Template.h"
 #include "llvm/ADT/STLForwardCompat.h"
@@ -3017,6 +3018,16 @@ static void checkNewAttributesAfterDef(Sema &S, Decl *New, const Decl *Old) {
       // declarations after definitions.
       ++I;
       continue;
+    } else if (isa<SYCLKernelEntryPointAttr>(NewAttribute)) {
+      // Elevate latent uses of the sycl_kernel_entry_point attribute to an
+      // error since the definition will have already been created without
+      // the semantic effects of the attribute having been applied.
+      S.Diag(NewAttribute->getLocation(),
+             diag::err_sycl_entry_point_after_definition);
+      S.Diag(Def->getLocation(), diag::note_previous_definition);
+      New->setInvalidDecl();
+      ++I;
+      continue;
     }
 
     S.Diag(NewAttribute->getLocation(),
@@ -12053,6 +12064,9 @@ bool Sema::CheckFunctionDeclaration(Scope *S, FunctionDecl *NewFD,
   if (LangOpts.OpenMP)
     OpenMP().ActOnFinishedFunctionDefinitionInOpenMPAssumeScope(NewFD);
 
+  if (LangOpts.isSYCL() && NewFD->hasAttr<SYCLKernelEntryPointAttr>())
+    SYCL().CheckSYCLEntryPointFunctionDecl(NewFD, /*CheckUseOfDecl=*/false);
+
   // Semantic checking for this function declaration (in isolation).
 
   if (getLangOpts().CPlusPlus) {
@@ -12285,6 +12299,13 @@ void Sema::CheckMain(FunctionDecl *FD, const DeclSpec &DS) {
     return;
   }
 
+  if (getLangOpts().isSYCL() && FD->hasAttr<SYCLKernelEntryPointAttr>()) {
+    Diag(FD->getAttr<SYCLKernelEntryPointAttr>()->getLocation(),
+         diag::err_sycl_entry_point_on_main);
+    FD->setInvalidDecl();
+    return;
+  }
+
   // Functions named main in hlsl are default entries, but don't have specific
   // signatures they are required to conform to.
   if (getLangOpts().HLSL)
@@ -15847,6 +15868,26 @@ Decl *Sema::ActOnFinishFunctionBody(Decl *dcl, Stmt *Body,
     // This is meant to pop the context added in ActOnStartOfFunctionDef().
     ExitFunctionBodyRAII ExitRAII(*this, isLambdaCallOperator(FD));
     if (FD) {
+      // Create SYCL kernel entry point function outline.
+      if (!FD->isInvalidDecl() && !FD->isDependentContext() &&
+          FD->hasAttr<SYCLKernelEntryPointAttr>()) {
+        if (FD->isDeleted()) {
+          Diag(FD->getAttr<SYCLKernelEntryPointAttr>()->getLocation(),
+              diag::err_sycl_entry_point_invalid)
+              << /*deleted function*/2;
+          FD->setInvalidDecl();
+        } else if (FD->isDefaulted()) {
+          Diag(FD->getAttr<SYCLKernelEntryPointAttr>()->getLocation(),
+              diag::err_sycl_entry_point_invalid)
+              << /*defaulted function*/3;
+          FD->setInvalidDecl();
+        } else if (FSI->isCoroutine()) {
+          Diag(FD->getAttr<SYCLKernelEntryPointAttr>()->getLocation(),
+              diag::err_sycl_entry_point_invalid)
+              << /*coroutine*/7;
+          FD->setInvalidDecl();
+        }
+      }
       // If this is called by Parser::ParseFunctionDefinition() after marking
       // the declaration as deleted, and if the deleted-function-body contains
       // a message (C++26), then a DefaultedOrDeletedInfo will have already been

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;