diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h
index b65a1f7dff5bc..c45f1ae77f17d 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,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
   //===--------------------------------------------------------------------===//
diff --git a/clang/include/clang/AST/SYCLKernelInfo.h b/clang/include/clang/AST/SYCLKernelInfo.h
new file mode 100644
index 0000000000000..79a83330f1d23
--- /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 ce86116680d7a..c4a3615752bf1 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 8ef151b3f2fdd..cd3aec8f70f02 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/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td
index 7d81bdf827ea0..2f4ca04e59620 100644
--- a/clang/include/clang/Basic/DiagnosticGroups.td
+++ b/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">;
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index ba813af960af6..a2bfedd6a993b 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/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>;
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index b86861ce7e8cf..877feaaaaf18d 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/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
diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h
index 27c42b5401830..5495349d151d5 100644
--- a/clang/include/clang/Sema/SemaSYCL.h
+++ b/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
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index 85b3984940ffc..1511bd8f6bf19 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/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();
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index de8805e15bc75..2e1c76326bd80 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/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
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 14cc51cf89665..2504fc2e7b8ec 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/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index 66df9c969256a..4d1b0ca7720ad 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -221,7 +221,8 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
                              ObjCInterfaceDecl *ClassReceiver,
                              bool SkipTrailingRequiresClause) {
   SourceLocation Loc = Locs.front();
-  if (getLangOpts().CPlusPlus && isa<FunctionDecl>(D)) {
+  if (FunctionDecl *FD = dyn_cast<FunctionDecl>(D);
+      getLangOpts().CPlusPlus && FD) {
     // If there were any diagnostics suppressed by template argument deduction,
     // emit them now.
     auto Pos = SuppressedDiagnostics.find(D->getCanonicalDecl());
@@ -236,6 +237,9 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
       Pos->second.clear();
     }
 
+    if (LangOpts.isSYCL() && FD->hasAttr<SYCLKernelEntryPointAttr>())
+      SYCL().CheckSYCLEntryPointFunctionDecl(FD, /*CheckUseOfDecl=*/true);
+
     // C++ [basic.start.main]p3:
     //   The function 'main' shall not be used within a program.
     if (cast<FunctionDecl>(D)->isMain())
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index f2d13d456c25f..6f8a3088d66b8 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -10,7 +10,9 @@
 
 #include "clang/Sema/SemaSYCL.h"
 #include "clang/AST/Mangle.h"
+#include "clang/AST/SYCLKernelInfo.h"
 #include "clang/AST/TypeOrdering.h"
+#include "clang/Basic/Diagnostic.h"
 #include "clang/Sema/Attr.h"
 #include "clang/Sema/ParsedAttr.h"
 #include "clang/Sema/Sema.h"
@@ -198,3 +200,168 @@ 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");
+
+  FunctionDecl *FD = dyn_cast<FunctionDecl>(D);
+  assert(FD && "Not a function decl");
+
+  bool hasError = false;
+  if (auto *MD = dyn_cast<CXXMethodDecl>(D)) {
+    if (!MD->isStatic()) {
+      Diag(AL.getLoc(), diag::err_sycl_entry_point_invalid)
+          << /*non-static member function*/ 0;
+      hasError = true;
+    }
+  }
+  if (FD->isVariadic()) {
+    Diag(AL.getLoc(), diag::err_sycl_entry_point_invalid)
+        << /*variadic function*/ 1;
+    hasError = true;
+  }
+  if (FD->isConsteval()) {
+    Diag(AL.getLoc(), diag::err_sycl_entry_point_invalid)
+        << /*consteval function*/ 5;
+    hasError = true;
+  } else if (FD->isConstexpr()) {
+    Diag(AL.getLoc(), diag::err_sycl_entry_point_invalid)
+        << /*constexpr function*/ 4;
+    hasError = true;
+  }
+  QualType Ret = FD->getReturnType();
+  if (!Ret->isDependentType() && !Ret->isVoidType()) {
+    Diag(AL.getLoc(), diag::err_sycl_entry_point_return_type);
+    hasError = true;
+  }
+
+  if (!hasError)
+    D->addAttr(::new (SemaRef.Context)
+                   SYCLKernelEntryPointAttr(SemaRef.Context, AL, TSI));
+}
+
+static SourceLocation SourceLocationForType(QualType QT) {
+  SourceLocation Loc;
+  const Type *T = QT->getUnqualifiedDesugaredType();
+  if (const TagType *TT = dyn_cast<TagType>(T))
+    Loc = TT->getDecl()->getLocation();
+  else if (const ObjCInterfaceType *ObjCIT = dyn_cast<ObjCInterfaceType>(T))
+    Loc = ObjCIT->getDecl()->getLocation();
+  return Loc;
+}
+
+static bool CheckSYCLKernelName(Sema &S, SourceLocation Loc,
+                                QualType KernelName) {
+  assert(!KernelName->isDependentType());
+
+  if (!KernelName->isStructureOrClassType()) {
+    // SYCL 2020 section 5.2, "Naming of kernels", only requires that the
+    // kernel name be a C++ typename. However, the definition of "kernel name"
+    // in the glossary states that a kernel name is a class type. Neither
+    // section explicitly states whether the kernel name type can be
+    // cv-qualified. For now, kernel name types are required to be class types
+    // and that they may be cv-qualified. The following issue requests
+    // clarification from the SYCL WG.
+    //   https://github.com/KhronosGroup/SYCL-Docs/issues/568
+    S.Diag(Loc, diag::warn_sycl_kernel_name_not_a_class_type)
+                  << KernelName;
+    SourceLocation DeclTypeLoc = SourceLocationForType(KernelName);
+    if (DeclTypeLoc.isValid())
+      S.Diag(DeclTypeLoc, diag::note_entity_declared_at) << KernelName;
+    return true;
+  }
+
+  return false;
+}
+
+void SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD, bool CheckUseOfDecl) {
+  // Delay diagnostic of templated decls until use or explicit instantiation.
+  //if (SemaRef.inTemplateInstantiation())
+  //  return;
+  if (FD->isTemplateInstantiation()) {
+    // For template instantiation, defer diagnostic until use.
+    if (!CheckUseOfDecl || DiagnosedSYCLKernelEntryPoint.count(FD))
+      return;
+    DiagnosedSYCLKernelEntryPoint.insert(FD);
+  }
+  // Ensure that all attributes present on the declaration are consistent
+  // and warn about any redundant ones.
+  const SYCLKernelEntryPointAttr *SKEPAttr = nullptr;
+  for (auto SAI = FD->specific_attr_begin<SYCLKernelEntryPointAttr>();
+       SAI != FD->specific_attr_end<SYCLKernelEntryPointAttr>(); ++SAI) {
+    if (!SKEPAttr) {
+      SKEPAttr = *SAI;
+      continue;
+    }
+    if (!getASTContext().hasSameType(SAI->getKernelName(),
+                                     SKEPAttr->getKernelName())) {
+      Diag(SAI->getLocation(), diag::err_sycl_entry_point_invalid_redeclaration)
+               << SAI->getKernelName() << SKEPAttr->getKernelName();
+      Diag(SKEPAttr->getLocation(), diag::note_previous_attribute);
+    } else {
+      Diag(SAI->getLocation(),
+           diag::warn_sycl_entry_point_redundant_declaration);
+      Diag(SKEPAttr->getLocation(), diag::note_previous_attribute);
+    }
+  }
+  assert(SKEPAttr && "Missing sycl_kernel_entry_point attribute");
+
+  // Ensure the kernel name type is valid.
+  if (!SKEPAttr->getKernelName()->isDependentType())
+    CheckSYCLKernelName(SemaRef, SKEPAttr->getLocation(),
+                        SKEPAttr->getKernelName());
+
+  // Ensure that an attribute present on the previous declaration
+  // matches the one on this declaration.
+  FunctionDecl *PrevFD = FD->getPreviousDecl();
+  if (PrevFD && !PrevFD->isInvalidDecl()) {
+    const auto *PrevSKEPAttr = PrevFD->getAttr<SYCLKernelEntryPointAttr>();
+    if (PrevSKEPAttr) {
+      if (!getASTContext().hasSameType(SKEPAttr->getKernelName(),
+                                       PrevSKEPAttr->getKernelName())) {
+        Diag(SKEPAttr->getLocation(),
+             diag::err_sycl_entry_point_invalid_redeclaration)
+                 << SKEPAttr->getKernelName()
+                 << PrevSKEPAttr->getKernelName();
+        Diag(PrevSKEPAttr->getLocation(),
+                                diag::note_previous_decl) << PrevFD;
+      }
+    }
+  }
+
+  if (FD->isNoReturn()) {
+    Diag(SKEPAttr->getLocation(), diag::err_sycl_entry_point_invalid)
+                                      << /*noreturn function*/ 6;
+  }
+
+  if (!FD->getReturnType()->isVoidType()) {
+    Diag(SKEPAttr->getLocation(),
+                            diag::err_sycl_entry_point_return_type);
+  }
+
+  if (!FD->isInvalidDecl() && !FD->isDependentContext()) {
+    const SYCLKernelInfo *SKI =
+        getASTContext().findSYCLKernelInfo(SKEPAttr->getKernelName());
+    if (SKI) {
+      if (!declaresSameEntity(FD, SKI->GetKernelEntryPointDecl())) {
+        Diag(FD->getLocation(), diag::err_sycl_kernel_name_conflict)
+                                    << SKEPAttr->getKernelName();
+        Diag(SKI->GetKernelEntryPointDecl()->getLocation(),
+             diag::note_previous_declaration);
+        for (const PartialDiagnosticAt &PD :
+             SYCLKernelEntryContextNotes.at(SKI->GetKernelEntryPointDecl()))
+          Diag(PD.first, PD.second);
+      }
+    } else {
+      // Note: In order to not interfere with SFINAE, we delay the diagnostic of
+      // conflicting names to when we act on the attribute.
+      ContextNotes &Notes = SYCLKernelEntryContextNotes[FD];
+      SemaRef.PrintInstantiationStack(
+          [&](const PartialDiagnosticAt &PD) { Notes.push_back(PD); });
+      getASTContext().registerSYCLEntryPointFunction(FD);
+    }
+  }
+}
diff --git a/clang/lib/Sema/SemaTemplate.cpp b/clang/lib/Sema/SemaTemplate.cpp
index 16f4542d78571..1df91c4f1e796 100644
--- a/clang/lib/Sema/SemaTemplate.cpp
+++ b/clang/lib/Sema/SemaTemplate.cpp
@@ -10364,6 +10364,9 @@ DeclResult Sema::ActOnExplicitInstantiation(Scope *S,
          diag::ext_explicit_instantiation_without_qualified_id)
     << Specialization << D.getCXXScopeSpec().getRange();
 
+  if (LangOpts.isSYCL() && Specialization->hasAttr<SYCLKernelEntryPointAttr>())
+    SYCL().CheckSYCLEntryPointFunctionDecl(Specialization, /*CheckUseOfDecl=*/true);
+
   CheckExplicitInstantiation(
       *this,
       FunTmpl ? (NamedDecl *)FunTmpl
diff --git a/clang/lib/Sema/SemaTemplateInstantiate.cpp b/clang/lib/Sema/SemaTemplateInstantiate.cpp
index 7481c700019dc..48ccc66016e84 100644
--- a/clang/lib/Sema/SemaTemplateInstantiate.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp
@@ -890,6 +890,15 @@ bool Sema::InstantiatingTemplate::CheckInstantiationDepth(
 }
 
 void Sema::PrintInstantiationStack() {
+  PrintInstantiationStack([&, this](const PartialDiagnosticAt &PD) {
+    DiagnosticBuilder Builder(Diags.Report(PD.first, PD.second.getDiagID()));
+    PD.second.Emit(Builder);
+  });
+}
+/// Prints the current instantiation stack through a series of
+/// notes.
+void Sema::PrintInstantiationStack(
+    std::function<void(const PartialDiagnosticAt &)> EmitDiag) {
   // Determine which template instantiations to skip, if any.
   unsigned SkipStart = CodeSynthesisContexts.size(), SkipEnd = SkipStart;
   unsigned Limit = Diags.getTemplateBacktraceLimit();
@@ -909,9 +918,9 @@ void Sema::PrintInstantiationStack() {
     if (InstantiationIdx >= SkipStart && InstantiationIdx < SkipEnd) {
       if (InstantiationIdx == SkipStart) {
         // Note that we're skipping instantiations.
-        Diags.Report(Active->PointOfInstantiation,
-                     diag::note_instantiation_contexts_suppressed)
-          << unsigned(CodeSynthesisContexts.size() - Limit);
+        EmitDiag({Active->PointOfInstantiation,
+                  PDiag(diag::note_instantiation_contexts_suppressed)
+                      << unsigned(CodeSynthesisContexts.size() - Limit)});
       }
       continue;
     }
@@ -923,37 +932,34 @@ void Sema::PrintInstantiationStack() {
         unsigned DiagID = diag::note_template_member_class_here;
         if (isa<ClassTemplateSpecializationDecl>(Record))
           DiagID = diag::note_template_class_instantiation_here;
-        Diags.Report(Active->PointOfInstantiation, DiagID)
-          << Record << Active->InstantiationRange;
+        EmitDiag({Active->PointOfInstantiation,
+                  PDiag(DiagID) << Record << Active->InstantiationRange});
       } else if (FunctionDecl *Function = dyn_cast<FunctionDecl>(D)) {
         unsigned DiagID;
         if (Function->getPrimaryTemplate())
           DiagID = diag::note_function_template_spec_here;
         else
           DiagID = diag::note_template_member_function_here;
-        Diags.Report(Active->PointOfInstantiation, DiagID)
-          << Function
-          << Active->InstantiationRange;
+        EmitDiag({Active->PointOfInstantiation,
+                  PDiag(DiagID) << Function << Active->InstantiationRange});
       } else if (VarDecl *VD = dyn_cast<VarDecl>(D)) {
-        Diags.Report(Active->PointOfInstantiation,
-                     VD->isStaticDataMember()?
-                       diag::note_template_static_data_member_def_here
-                     : diag::note_template_variable_def_here)
-          << VD
-          << Active->InstantiationRange;
+        EmitDiag({Active->PointOfInstantiation,
+                  PDiag(VD->isStaticDataMember()
+                            ? diag::note_template_static_data_member_def_here
+                            : diag::note_template_variable_def_here)
+                      << VD << Active->InstantiationRange});
       } else if (EnumDecl *ED = dyn_cast<EnumDecl>(D)) {
-        Diags.Report(Active->PointOfInstantiation,
-                     diag::note_template_enum_def_here)
-          << ED
-          << Active->InstantiationRange;
+        EmitDiag({Active->PointOfInstantiation,
+                  PDiag(diag::note_template_enum_def_here)
+                      << ED << Active->InstantiationRange});
       } else if (FieldDecl *FD = dyn_cast<FieldDecl>(D)) {
-        Diags.Report(Active->PointOfInstantiation,
-                     diag::note_template_nsdmi_here)
-            << FD << Active->InstantiationRange;
+        EmitDiag({Active->PointOfInstantiation,
+                  PDiag(diag::note_template_nsdmi_here)
+                      << FD << Active->InstantiationRange});
       } else if (ClassTemplateDecl *CTD = dyn_cast<ClassTemplateDecl>(D)) {
-        Diags.Report(Active->PointOfInstantiation,
-                     diag::note_template_class_instantiation_here)
-            << CTD << Active->InstantiationRange;
+        EmitDiag({Active->PointOfInstantiation,
+                  PDiag(diag::note_template_class_instantiation_here)
+                      << CTD << Active->InstantiationRange});
       }
       break;
     }
@@ -965,35 +971,35 @@ void Sema::PrintInstantiationStack() {
       Template->printName(OS, getPrintingPolicy());
       printTemplateArgumentList(OS, Active->template_arguments(),
                                 getPrintingPolicy());
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_default_arg_instantiation_here)
-        << OS.str()
-        << Active->InstantiationRange;
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_default_arg_instantiation_here)
+                    << OS.str() << Active->InstantiationRange});
       break;
     }
 
     case CodeSynthesisContext::ExplicitTemplateArgumentSubstitution: {
       FunctionTemplateDecl *FnTmpl = cast<FunctionTemplateDecl>(Active->Entity);
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_explicit_template_arg_substitution_here)
-        << FnTmpl
-        << getTemplateArgumentBindingsText(FnTmpl->getTemplateParameters(),
-                                           Active->TemplateArgs,
-                                           Active->NumTemplateArgs)
-        << Active->InstantiationRange;
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_explicit_template_arg_substitution_here)
+                    << FnTmpl
+                    << getTemplateArgumentBindingsText(
+                           FnTmpl->getTemplateParameters(),
+                           Active->TemplateArgs, Active->NumTemplateArgs)
+                    << Active->InstantiationRange});
       break;
     }
 
     case CodeSynthesisContext::DeducedTemplateArgumentSubstitution: {
       if (FunctionTemplateDecl *FnTmpl =
               dyn_cast<FunctionTemplateDecl>(Active->Entity)) {
-        Diags.Report(Active->PointOfInstantiation,
-                     diag::note_function_template_deduction_instantiation_here)
-          << FnTmpl
-          << getTemplateArgumentBindingsText(FnTmpl->getTemplateParameters(),
-                                             Active->TemplateArgs,
-                                             Active->NumTemplateArgs)
-          << Active->InstantiationRange;
+        EmitDiag(
+            {Active->PointOfInstantiation,
+             PDiag(diag::note_function_template_deduction_instantiation_here)
+                 << FnTmpl
+                 << getTemplateArgumentBindingsText(
+                        FnTmpl->getTemplateParameters(), Active->TemplateArgs,
+                        Active->NumTemplateArgs)
+                 << Active->InstantiationRange});
       } else {
         bool IsVar = isa<VarTemplateDecl>(Active->Entity) ||
                      isa<VarTemplateSpecializationDecl>(Active->Entity);
@@ -1012,12 +1018,13 @@ void Sema::PrintInstantiationStack() {
           llvm_unreachable("unexpected template kind");
         }
 
-        Diags.Report(Active->PointOfInstantiation,
-                     diag::note_deduced_template_arg_substitution_here)
-          << IsVar << IsTemplate << cast<NamedDecl>(Active->Entity)
-          << getTemplateArgumentBindingsText(Params, Active->TemplateArgs,
-                                             Active->NumTemplateArgs)
-          << Active->InstantiationRange;
+        EmitDiag(
+            {Active->PointOfInstantiation,
+             PDiag(diag::note_deduced_template_arg_substitution_here)
+                 << IsVar << IsTemplate << cast<NamedDecl>(Active->Entity)
+                 << getTemplateArgumentBindingsText(
+                        Params, Active->TemplateArgs, Active->NumTemplateArgs)
+                 << Active->InstantiationRange});
       }
       break;
     }
@@ -1031,10 +1038,9 @@ void Sema::PrintInstantiationStack() {
       FD->printName(OS, getPrintingPolicy());
       printTemplateArgumentList(OS, Active->template_arguments(),
                                 getPrintingPolicy());
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_default_function_arg_instantiation_here)
-        << OS.str()
-        << Active->InstantiationRange;
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_default_function_arg_instantiation_here)
+                    << OS.str() << Active->InstantiationRange});
       break;
     }
 
@@ -1051,14 +1057,13 @@ void Sema::PrintInstantiationStack() {
         TemplateParams =
           cast<ClassTemplatePartialSpecializationDecl>(Active->Template)
                                                       ->getTemplateParameters();
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_prior_template_arg_substitution)
-        << isa<TemplateTemplateParmDecl>(Parm)
-        << Name
-        << getTemplateArgumentBindingsText(TemplateParams,
-                                           Active->TemplateArgs,
-                                           Active->NumTemplateArgs)
-        << Active->InstantiationRange;
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_prior_template_arg_substitution)
+                    << isa<TemplateTemplateParmDecl>(Parm) << Name
+                    << getTemplateArgumentBindingsText(TemplateParams,
+                                                       Active->TemplateArgs,
+                                                       Active->NumTemplateArgs)
+                    << Active->InstantiationRange});
       break;
     }
 
@@ -1071,55 +1076,56 @@ void Sema::PrintInstantiationStack() {
           cast<ClassTemplatePartialSpecializationDecl>(Active->Template)
                                                       ->getTemplateParameters();
 
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_template_default_arg_checking)
-        << getTemplateArgumentBindingsText(TemplateParams,
-                                           Active->TemplateArgs,
-                                           Active->NumTemplateArgs)
-        << Active->InstantiationRange;
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_template_default_arg_checking)
+                    << getTemplateArgumentBindingsText(TemplateParams,
+                                                       Active->TemplateArgs,
+                                                       Active->NumTemplateArgs)
+                    << Active->InstantiationRange});
       break;
     }
 
     case CodeSynthesisContext::ExceptionSpecEvaluation:
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_evaluating_exception_spec_here)
-          << cast<FunctionDecl>(Active->Entity);
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_evaluating_exception_spec_here)
+                    << cast<FunctionDecl>(Active->Entity)});
       break;
 
     case CodeSynthesisContext::ExceptionSpecInstantiation:
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_template_exception_spec_instantiation_here)
-        << cast<FunctionDecl>(Active->Entity)
-        << Active->InstantiationRange;
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_template_exception_spec_instantiation_here)
+                    << cast<FunctionDecl>(Active->Entity)
+                    << Active->InstantiationRange});
       break;
 
     case CodeSynthesisContext::RequirementInstantiation:
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_template_requirement_instantiation_here)
-        << Active->InstantiationRange;
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_template_requirement_instantiation_here)
+                    << Active->InstantiationRange});
       break;
     case CodeSynthesisContext::RequirementParameterInstantiation:
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_template_requirement_params_instantiation_here)
-          << Active->InstantiationRange;
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_template_requirement_params_instantiation_here)
+                    << Active->InstantiationRange});
       break;
 
     case CodeSynthesisContext::NestedRequirementConstraintsCheck:
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_nested_requirement_here)
-        << Active->InstantiationRange;
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_nested_requirement_here)
+                    << Active->InstantiationRange});
       break;
 
     case CodeSynthesisContext::DeclaringSpecialMember:
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_in_declaration_of_implicit_special_member)
-          << cast<CXXRecordDecl>(Active->Entity)
-          << llvm::to_underlying(Active->SpecialMember);
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_in_declaration_of_implicit_special_member)
+                    << cast<CXXRecordDecl>(Active->Entity)
+                    << llvm::to_underlying(Active->SpecialMember)});
       break;
 
     case CodeSynthesisContext::DeclaringImplicitEqualityComparison:
-      Diags.Report(Active->Entity->getLocation(),
-                   diag::note_in_declaration_of_implicit_equality_comparison);
+      EmitDiag(
+          {Active->Entity->getLocation(),
+           PDiag(diag::note_in_declaration_of_implicit_equality_comparison)});
       break;
 
     case CodeSynthesisContext::DefiningSynthesizedFunction: {
@@ -1130,60 +1136,62 @@ void Sema::PrintInstantiationStack() {
           FD ? getDefaultedFunctionKind(FD) : DefaultedFunctionKind();
       if (DFK.isSpecialMember()) {
         auto *MD = cast<CXXMethodDecl>(FD);
-        Diags.Report(Active->PointOfInstantiation,
-                     diag::note_member_synthesized_at)
-            << MD->isExplicitlyDefaulted()
-            << llvm::to_underlying(DFK.asSpecialMember())
-            << Context.getTagDeclType(MD->getParent());
+        EmitDiag({Active->PointOfInstantiation,
+                  PDiag(diag::note_member_synthesized_at)
+                      << MD->isExplicitlyDefaulted()
+                      << llvm::to_underlying(DFK.asSpecialMember())
+                      << Context.getTagDeclType(MD->getParent())});
       } else if (DFK.isComparison()) {
         QualType RecordType = FD->getParamDecl(0)
                                   ->getType()
                                   .getNonReferenceType()
                                   .getUnqualifiedType();
-        Diags.Report(Active->PointOfInstantiation,
-                     diag::note_comparison_synthesized_at)
-            << (int)DFK.asComparison() << RecordType;
+        EmitDiag({Active->PointOfInstantiation,
+                  PDiag(diag::note_comparison_synthesized_at)
+                      << (int)DFK.asComparison() << RecordType});
       }
       break;
     }
 
     case CodeSynthesisContext::RewritingOperatorAsSpaceship:
-      Diags.Report(Active->Entity->getLocation(),
-                   diag::note_rewriting_operator_as_spaceship);
+      EmitDiag({Active->Entity->getLocation(),
+                PDiag(diag::note_rewriting_operator_as_spaceship)});
       break;
 
     case CodeSynthesisContext::InitializingStructuredBinding:
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_in_binding_decl_init)
-          << cast<BindingDecl>(Active->Entity);
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_in_binding_decl_init)
+                    << cast<BindingDecl>(Active->Entity)});
       break;
 
     case CodeSynthesisContext::MarkingClassDllexported:
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_due_to_dllexported_class)
-          << cast<CXXRecordDecl>(Active->Entity) << !getLangOpts().CPlusPlus11;
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_due_to_dllexported_class)
+                    << cast<CXXRecordDecl>(Active->Entity)
+                    << !getLangOpts().CPlusPlus11});
       break;
 
     case CodeSynthesisContext::BuildingBuiltinDumpStructCall:
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_building_builtin_dump_struct_call)
-          << convertCallArgsToString(
-                 *this, llvm::ArrayRef(Active->CallArgs, Active->NumCallArgs));
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_building_builtin_dump_struct_call)
+                    << convertCallArgsToString(
+                           *this, llvm::ArrayRef(Active->CallArgs,
+                                                 Active->NumCallArgs))});
       break;
 
     case CodeSynthesisContext::Memoization:
       break;
 
     case CodeSynthesisContext::LambdaExpressionSubstitution:
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_lambda_substitution_here);
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_lambda_substitution_here)});
       break;
     case CodeSynthesisContext::ConstraintsCheck: {
       unsigned DiagID = 0;
       if (!Active->Entity) {
-        Diags.Report(Active->PointOfInstantiation,
-                     diag::note_nested_requirement_here)
-          << Active->InstantiationRange;
+        EmitDiag({Active->PointOfInstantiation,
+                  PDiag(diag::note_nested_requirement_here)
+                      << Active->InstantiationRange});
         break;
       }
       if (isa<ConceptDecl>(Active->Entity))
@@ -1205,34 +1213,34 @@ void Sema::PrintInstantiationStack() {
         printTemplateArgumentList(OS, Active->template_arguments(),
                                   getPrintingPolicy());
       }
-      Diags.Report(Active->PointOfInstantiation, DiagID) << OS.str()
-        << Active->InstantiationRange;
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(DiagID) << OS.str() << Active->InstantiationRange});
       break;
     }
     case CodeSynthesisContext::ConstraintSubstitution:
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_constraint_substitution_here)
-          << Active->InstantiationRange;
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_constraint_substitution_here)
+                    << Active->InstantiationRange});
       break;
     case CodeSynthesisContext::ConstraintNormalization:
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_constraint_normalization_here)
-          << cast<NamedDecl>(Active->Entity) << Active->InstantiationRange;
+      EmitDiag({Active->PointOfInstantiation,
+                   PDiag(diag::note_constraint_normalization_here)
+          << cast<NamedDecl>(Active->Entity) << Active->InstantiationRange});
       break;
     case CodeSynthesisContext::ParameterMappingSubstitution:
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_parameter_mapping_substitution_here)
-          << Active->InstantiationRange;
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_parameter_mapping_substitution_here)
+                    << Active->InstantiationRange});
       break;
     case CodeSynthesisContext::BuildingDeductionGuides:
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_building_deduction_guide_here);
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_building_deduction_guide_here)});
       break;
     case CodeSynthesisContext::TypeAliasTemplateInstantiation:
-      Diags.Report(Active->PointOfInstantiation,
-                   diag::note_template_type_alias_instantiation_here)
-          << cast<TypeAliasTemplateDecl>(Active->Entity)
-          << Active->InstantiationRange;
+      EmitDiag({Active->PointOfInstantiation,
+                PDiag(diag::note_template_type_alias_instantiation_here)
+                    << cast<TypeAliasTemplateDecl>(Active->Entity)
+                    << Active->InstantiationRange});
       break;
     }
   }
diff --git a/clang/lib/Serialization/ASTReaderDecl.cpp b/clang/lib/Serialization/ASTReaderDecl.cpp
index 321e65fd2b094..02e723aef0ff8 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 0000000000000..ab76cfe9a2c95
--- /dev/null
+++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
@@ -0,0 +1,154 @@
+// 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>
+
+// 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 a diagnostic
+// FIXME: being incorrectly issued for the following explicit specialization
+// FIXME: due to conflicting kernel name types (KN<5,3> vs the incorrectly
+// FIXME: inherited KN<5,-1>).
+#if 0
+template<>
+[[clang::sycl_kernel_entry_point(KN<5,3>)]]
+void skep5<KN<5,-1>>(long long) {
+}
+// FIXME-CHECK: |-FunctionDecl {{.*}} prev {{.*}} skep5 'void (long long)' explicit_specialization
+// FIXME-CHECK: | |-TemplateArgument type 'KN<5, -1>'
+// FIXME-CHECK: | |-TemplateArgument type 'long long'
+// FIXME-CHECK: | `-SYCLKernelEntryPointAttr {{.*}} KN<5, 3>
+#endif
+
+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 914f94c08a9fd..5c2f3a347dfb7 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-appertainment.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp
new file mode 100644
index 0000000000000..b70f7d1254a24
--- /dev/null
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-appertainment.cpp
@@ -0,0 +1,239 @@
+// 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 appertainment for the sycl_kernel_entry_point attribute.
+
+#if __cplusplus >= 202002L
+// Mock coroutine support.
+namespace std {
+
+template<typename Promise = void>
+struct coroutine_handle {
+  template<typename T>
+  coroutine_handle(const coroutine_handle<T>&);
+  static coroutine_handle from_address(void *addr);
+};
+
+template<typename R, typename... Args>
+struct coroutine_traits {
+  struct suspend_never {
+    bool await_ready() const noexcept;
+    void await_suspend(std::coroutine_handle<>) const noexcept;
+    void await_resume() const noexcept;
+  };
+  struct promise_type {
+    void get_return_object() noexcept;
+    suspend_never initial_suspend() const noexcept;
+    suspend_never final_suspend() const noexcept;
+    void return_void() noexcept;
+    void unhandled_exception() noexcept;
+  };
+};
+
+}
+#endif
+
+// A unique kernel name type is required for each declared kernel entry point.
+template<int, int = 0> struct KN;
+
+
+////////////////////////////////////////////////////////////////////////////////
+// Valid declarations.
+////////////////////////////////////////////////////////////////////////////////
+
+// Function declaration with GNU attribute spelling
+__attribute__((sycl_kernel_entry_point(KN<1>)))
+void ok1();
+
+// Function declaration with Clang attribute spelling.
+[[clang::sycl_kernel_entry_point(KN<2>)]]
+void ok2();
+
+// Function definition.
+[[clang::sycl_kernel_entry_point(KN<3>)]]
+void ok3() {}
+
+// Function template definition.
+template<typename KNT, typename T>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void ok4(T) {}
+
+// Function template explicit specialization.
+template<>
+[[clang::sycl_kernel_entry_point(KN<4,1>)]]
+void ok4<KN<4,1>>(int) {}
+
+// Function template explicit instantiation.
+template void ok4<KN<4,2>, long>(long);
+
+namespace NS {
+// Function declaration at namespace scope.
+[[clang::sycl_kernel_entry_point(KN<5>)]]
+void ok5();
+}
+
+struct S6 {
+  // Static member function declaration.
+  [[clang::sycl_kernel_entry_point(KN<6>)]]
+  static void ok6();
+};
+
+// The sycl_kernel_entry_point attribute must match across declarations and
+// cannot be added for the first time after a definition.
+[[clang::sycl_kernel_entry_point(KN<7>)]]
+void ok7();
+[[clang::sycl_kernel_entry_point(KN<7>)]]
+void ok7();
+[[clang::sycl_kernel_entry_point(KN<8>)]]
+void ok8();
+void ok8() {}
+void ok9();
+[[clang::sycl_kernel_entry_point(KN<9>)]]
+void ok9() {}
+
+using VOID = void;
+[[clang::sycl_kernel_entry_point(KN<10>)]]
+VOID ok10();
+[[clang::sycl_kernel_entry_point(KN<11>)]]
+const void ok11();
+
+
+////////////////////////////////////////////////////////////////////////////////
+// Invalid declarations.
+////////////////////////////////////////////////////////////////////////////////
+
+struct Smain;
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute only applies to functions with a 'void' return type}}
+[[clang::sycl_kernel_entry_point(Smain)]]
+int main();
+
+template<int> struct BADKN;
+
+struct B1 {
+  // Non-static data member declaration.
+  // expected-error@+1 {{'sycl_kernel_entry_point' attribute only applies to functions}}
+  [[clang::sycl_kernel_entry_point(BADKN<1>)]]
+  int bad1;
+};
+
+struct B2 {
+  // Static data member declaration.
+  // expected-error@+1 {{'sycl_kernel_entry_point' attribute only applies to functions}}
+  [[clang::sycl_kernel_entry_point(BADKN<2>)]]
+  static int bad2;
+};
+
+struct B3 {
+  // Non-static member function declaration.
+  // expected-error@+1 {{'sycl_kernel_entry_point' attribute cannot be applied to a non-static member function}}
+  [[clang::sycl_kernel_entry_point(BADKN<3>)]]
+  void bad3();
+};
+
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute only applies to functions}}
+namespace bad4 [[clang::sycl_kernel_entry_point(BADKN<4>)]] {}
+
+#if __cplusplus >= 202002L
+// expected-error@+2 {{'sycl_kernel_entry_point' attribute only applies to functions}}
+template<typename>
+concept bad5 [[clang::sycl_kernel_entry_point(BADKN<5>)]] = true;
+#endif
+
+// Type alias declarations.
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute only applies to functions}}
+typedef void bad6 [[clang::sycl_kernel_entry_point(BADKN<6>)]] ();
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute only applies to functions}}
+using bad7 [[clang::sycl_kernel_entry_point(BADKN<7>)]] = void();
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute only applies to functions}}
+using bad8 [[clang::sycl_kernel_entry_point(BADKN<8>)]] = int;
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute cannot be applied to types}}
+using bad9 = int [[clang::sycl_kernel_entry_point(BADKN<9>)]];
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute cannot be applied to types}}
+using bad10 = int() [[clang::sycl_kernel_entry_point(BADKN<10>)]];
+
+// Variable declaration.
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute only applies to functions}}
+[[clang::sycl_kernel_entry_point(BADKN<11>)]]
+int bad11;
+
+// Class declaration.
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute only applies to functions}}
+struct [[clang::sycl_kernel_entry_point(BADKN<12>)]] bad12;
+
+// Enumeration declaration.
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute only applies to functions}}
+enum [[clang::sycl_kernel_entry_point(BADKN<13>)]] bad13 {};
+
+// Enumerator.
+// expected-error@+2 {{'sycl_kernel_entry_point' attribute only applies to functions}}
+enum {
+  bad14 [[clang::sycl_kernel_entry_point(BADKN<14>)]]
+};
+
+// Attribute added after the definition.
+// expected-error@+3 {{'sycl_kernel_entry_point' attribute cannot be added to a function after the function is defined}}
+// expected-note@+1 {{previous definition is here}}
+void bad15() {}
+[[clang::sycl_kernel_entry_point(BADKN<15>)]]
+void bad15();
+
+// The function must return void.
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute only applies to functions with a 'void' return type}}
+[[clang::sycl_kernel_entry_point(BADKN<16>)]]
+int bad16();
+
+// Function parameters.
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute only applies to functions}}
+void bad17(void (fp [[clang::sycl_kernel_entry_point(BADKN<17>)]])());
+
+// Function template parameters.
+// FIXME-expected-error@+1 {{'sycl_kernel_entry_point' attribute only applies to functions}}
+template<void (fp [[clang::sycl_kernel_entry_point(BADKN<18>)]])()>
+void bad18();
+
+#if __cplusplus >= 202002L
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute cannot be applied to a coroutine function}}
+[[clang::sycl_kernel_entry_point(BADKN<19>)]]
+void bad19() {
+  co_return;
+}
+#endif
+
+struct B20 {
+  // expected-error@+1 {{'sycl_kernel_entry_point' attribute cannot be applied to a non-static member function}}
+  [[clang::sycl_kernel_entry_point(BADKN<20>)]]
+  B20();
+};
+
+struct B21 {
+  // expected-error@+1 {{'sycl_kernel_entry_point' attribute cannot be applied to a non-static member function}}
+  [[clang::sycl_kernel_entry_point(BADKN<21>)]]
+  ~B21();
+};
+
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute cannot be applied to a variadic function}}
+[[clang::sycl_kernel_entry_point(BADKN<22>)]]
+void bad22(...);
+
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute cannot be applied to a deleted function}}
+[[clang::sycl_kernel_entry_point(BADKN<23>)]]
+void bad23() = delete;
+
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute cannot be applied to a constexpr function}}
+[[clang::sycl_kernel_entry_point(BADKN<24>)]]
+constexpr void bad24() {}
+
+#if __cplusplus >= 202002L
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute cannot be applied to a consteval function}}
+[[clang::sycl_kernel_entry_point(BADKN<25>)]]
+consteval void bad25() {}
+#endif
+
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute cannot be applied to a noreturn function}}
+[[clang::sycl_kernel_entry_point(BADKN<26>)]]
+[[noreturn]] void bad26();
+
+// expected-error@+3 {{attribute 'target' multiversioning cannot be combined with attribute 'sycl_kernel_entry_point'}}
+__attribute__((target("avx"))) void bad27();
+[[clang::sycl_kernel_entry_point(BADKN<27>)]]
+__attribute__((target("sse4.2"))) void bad27();
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-argument.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-argument.cpp
new file mode 100644
index 0000000000000..a745a5da8889f
--- /dev/null
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-argument.cpp
@@ -0,0 +1,268 @@
+// 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 proper handling of a sycl_kernel_entry_point attribute
+// argument list. The single argument is required to denote a class type that
+// meets the requirements of a SYCL kernel name as described in section 5.2,
+// "Naming of kernels", of the SYCL 2020 specification.
+
+// Common entities used to validate kernel name arguments.
+template<int> struct ST; // #ST-decl
+template<int N> using TTA = ST<N>; // #TTA-decl
+
+
+////////////////////////////////////////////////////////////////////////////////
+// Valid declarations.
+////////////////////////////////////////////////////////////////////////////////
+struct S1;
+[[clang::sycl_kernel_entry_point(S1)]] void ok1();
+
+typedef struct {} TA2;
+[[clang::sycl_kernel_entry_point(TA2)]] void ok2();
+
+using TA3 = struct {};
+[[clang::sycl_kernel_entry_point(TA3)]] void ok3();
+
+[[clang::sycl_kernel_entry_point(ST<4>)]] void ok4();
+
+[[clang::sycl_kernel_entry_point(TTA<5>)]] void ok5();
+
+namespace NS6 {
+  struct NSS;
+}
+[[clang::sycl_kernel_entry_point(NS6::NSS)]] void ok6();
+
+namespace {
+  struct UNSS7;
+}
+[[clang::sycl_kernel_entry_point(UNSS7)]] void ok7();
+
+struct {} s;
+[[clang::sycl_kernel_entry_point(decltype(s))]] void ok8();
+
+template<typename KN>
+[[clang::sycl_kernel_entry_point(KN)]] void ok9();
+void test_ok9() {
+  ok9<struct LS1>();
+}
+
+template<int, typename KN>
+[[clang::sycl_kernel_entry_point(KN)]] void ok10();
+void test_ok10() {
+  ok10<1, struct LS2>();
+}
+
+namespace NS11 {
+  struct NSS;
+}
+template<typename T>
+[[clang::sycl_kernel_entry_point(T)]] void ok11() {}
+template<>
+[[clang::sycl_kernel_entry_point(NS11::NSS)]] void ok11<NS11::NSS>() {}
+
+struct S12;
+[[clang::sycl_kernel_entry_point(S12)]] void ok12();
+[[clang::sycl_kernel_entry_point(S12)]] void ok12() {}
+
+struct S13;
+struct S13;
+// expected-warning@+3 {{redundant 'sycl_kernel_entry_point' attribute}}
+// expected-note@+1  {{previous attribute is here}}
+[[clang::sycl_kernel_entry_point(S13),
+  clang::sycl_kernel_entry_point(S13)]]
+void ok13();
+
+template<typename T>
+[[clang::sycl_kernel_entry_point(T)]] void ok14(T k);
+void test_ok14() {
+  ok14([]{});
+}
+
+
+////////////////////////////////////////////////////////////////////////////////
+// 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();
+
+// expected-error@+1 {{'int' is not a valid SYCL kernel name type; a class type is required}}
+[[clang::sycl_kernel_entry_point(int)]] void bad6();
+
+// expected-error@+1 {{'int ()' is not a valid SYCL kernel name type; a class type is required}}
+[[clang::sycl_kernel_entry_point(int())]] void bad7();
+
+// expected-error@+1 {{'int (*)()' is not a valid SYCL kernel name type; a class type is required}}
+[[clang::sycl_kernel_entry_point(int(*)())]] void bad8();
+
+// expected-error@+1 {{'int (&)()' is not a valid SYCL kernel name type; a class type is required}}
+[[clang::sycl_kernel_entry_point(int(&)())]] void bad9();
+
+// expected-error@+1 {{'decltype(nullptr)' (aka 'std::nullptr_t') is not a valid SYCL kernel name type; a class type is required}}
+[[clang::sycl_kernel_entry_point(decltype(nullptr))]] void bad10();
+
+void f11();
+// expected-error@+1 {{unknown type name 'f11'}}
+[[clang::sycl_kernel_entry_point(f11)]] void bad11();
+
+// 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 bad12();
+
+// 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 bad13();
+
+union U; // #U-decl
+// expected-error@+2 {{'U' is not a valid SYCL kernel name type; a class type is required}}
+// expected-note@#U-decl {{'U' declared here}}
+[[clang::sycl_kernel_entry_point(U)]] void bad14();
+
+enum E15 {}; // #E15-decl
+// expected-error@+2 {{'E15' is not a valid SYCL kernel name type; a class type is required}}
+// expected-note@#E15-decl {{'E15' declared here}}
+[[clang::sycl_kernel_entry_point(E15)]] void bad15();
+
+enum E16 : int; // #E16-decl
+// expected-error@+2 {{'E16' is not a valid SYCL kernel name type; a class type is required}}
+// expected-note@#E16-decl {{'E16' declared here}}
+[[clang::sycl_kernel_entry_point(E16)]] void bad16();
+
+enum {
+  e17
+};
+// expected-error@+1 {{unknown type name 'e17'}}
+[[clang::sycl_kernel_entry_point(e17)]] void bad17();
+
+#if __cplusplus >= 202002L
+template<typename> concept C = true;
+// expected-error@+1 {{expected a type}}
+[[clang::sycl_kernel_entry_point(C)]] void bad18();
+
+// expected-error@+1 {{expected a type}}
+[[clang::sycl_kernel_entry_point(C<int>)]] void bad19();
+#endif
+
+struct B20 {
+  struct MS;
+};
+// FIXME-expected-error@+1 {{'sycl_kernel_entry_point' attribute argument must be a forward declarable class type}}
+[[clang::sycl_kernel_entry_point(B20::MS)]] void bad20();
+
+struct B21 {
+  struct MS;
+};
+// FIXME-expected-error@+3 {{'sycl_kernel_entry_point' attribute argument must be a forward declarable class type}}
+template<typename T>
+[[clang::sycl_kernel_entry_point(typename T::MS)]] void bad21() {}
+template void bad21<B21>();
+
+struct B22; // #B22-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 'B22' named in nested name specifier}}
+// expected-note@+5 {{in instantiation of function template specialization 'bad22<B22>' requested here}}
+// expected-note@#B22-decl {{forward declaration of 'B22'}}
+template<typename T>
+[[clang::sycl_kernel_entry_point(typename T::not_found)]] void bad22() {}
+template<>
+void bad22<B22>() {}
+
+template<typename T>
+[[clang::sycl_kernel_entry_point(T)]] void bad23();
+void f23() {
+  // FIXME-expected-error@+2 {{'sycl_kernel_entry_point' attribute argument must be a forward declarable class type}}
+  struct LS;
+  bad23<LS>();
+}
+
+struct B24_1;
+struct B24_2;
+// expected-error@+3 {{'sycl_kernel_entry_point' kernel name argument does not match prior declaration: 'B24_2' vs 'B24_1'}}
+// expected-note@+1  {{'bad24' declared here}}
+[[clang::sycl_kernel_entry_point(B24_1)]] void bad24();
+[[clang::sycl_kernel_entry_point(B24_2)]] void bad24() {}
+
+struct B25_1;
+struct B25_2;
+// expected-error@+3 {{'sycl_kernel_entry_point' kernel name argument does not match prior declaration: 'B25_2' vs 'B25_1'}}
+// expected-note@+1  {{previous attribute is here}}
+[[clang::sycl_kernel_entry_point(B25_1),
+  clang::sycl_kernel_entry_point(B25_2)]]
+void bad25();
+
+// Validate that conflicting kernel names are diagnosed for non-defining declarations.
+struct B26;
+// expected-error@+3 {{'sycl_kernel_entry_point' kernel name 'B26' conflicts with a previous declaration}}
+// expected-note@+1  {{previous declaration is here}}
+[[clang::sycl_kernel_entry_point(B26)]] void bad26_1();
+[[clang::sycl_kernel_entry_point(B26)]] void bad26_2();
+
+
+struct B27;
+
+template<typename KN, typename K>
+[[clang::sycl_kernel_entry_point(KN)]]
+void bad27(K ker, int i) { ker(); } // #bad27-decl
+
+// Overload to ensure attribute checks don't affect overload resolution.
+template<typename KN, typename K>
+void bad27(K ker, long i) { ker(); }
+
+void test_bad27_1() {
+  int i = 0;
+  long l = 0;
+  // expected-error@#bad27-decl {{'sycl_kernel_entry_point' kernel name 'B27' conflicts with a previous declaration}}
+  // expected-note-re@+4  {{in instantiation of function template specialization 'bad27<B27, (lambda at{{.*}}}}
+  // expected-note@#bad27-decl  {{previous declaration is here}}
+  // expected-note-re@+1  {{in instantiation of function template specialization 'bad27<B27, (lambda at{{.*}}}}
+  bad27<B27>([]{}, i); // #bad27-FirstUse
+  bad27<B27>([]{}, i);
+  // this doesn't trigger any error as overload doesn't have the attribute.
+  bad27<B27>([]{}, l);
+}
+
+void test_bad27_2() {
+  int i = 0;
+  // expected-error@#bad27-decl {{'sycl_kernel_entry_point' kernel name 'B27' conflicts with a previous declaration}}
+  // expected-note-re@+3  {{in instantiation of function template specialization 'bad27<B27, (lambda at{{.*}}}}
+  // expected-note@#bad27-decl  {{previous declaration is here}}
+  // expected-note-re@#bad27-FirstUse  {{in instantiation of function template specialization 'bad27<B27, (lambda at{{.*}}}}
+  bad27<B27>([]{}, i);
+}
+
+template<typename T>
+void test_bad27_3() {
+  int i = 0;
+  // expected-error@#bad27-decl {{'sycl_kernel_entry_point' kernel name 'B27' conflicts with a previous declaration}}
+  // expected-note-re@+4  {{in instantiation of function template specialization 'bad27<B27, (lambda at{{.*}}}}
+  // expected-note@#test_bad27_3-call  {{in instantiation of function template specialization 'test_bad27_3<B27>' requested here}}
+  // expected-note@#bad27-decl  {{previous declaration is here}}
+  // expected-note-re@#bad27-FirstUse  {{in instantiation of function template specialization 'bad27<B27, (lambda at{{.*}}}}
+  bad27<T>([]{}, i);
+}
+
+void test_bad27_4() {
+  test_bad27_3<B27>(); // #test_bad27_3-call
+}
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 0000000000000..30de6ae0b0e6f
--- /dev/null
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-ignored.cpp
@@ -0,0 +1,17 @@
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -verify %s
+
+// These tests validate that the sycl_kernel_entry_point attribute is ignored
+// when SYCL support is not enabled.
+
+// A unique kernel name type is required for each declared kernel entry point.
+template<int> struct KN;
+
+// expected-warning@+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<typename KNT>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void ok2() {}
+template void ok2<KN<2>>();