From 3c4a2b8a52d3f1c730df88a308dece21a67834ef Mon Sep 17 00:00:00 2001
From: Tom Honermann <tom.honermann@intel.com>
Date: Fri, 4 Oct 2024 11:10:32 -0700
Subject: [PATCH 1/9] [SYCL] The sycl_kernel_entry_point attribute.

The `sycl_kernel_entry_point` attribute is used to declare a function that
defines a pattern for an offload kernel to be emitted. The attribute requires
a single type argument that specifies the type used as a SYCL kernel name as
described in section 5.2, "Naming of kernels", of the SYCL 2020 specification.

Properties of the offload kernel are collected when a function declared with
the `sycl_kernel_entry_point` attribute is parsed or instantiated. These
properties, such as the kernel name type, are stored in the AST context where
they are (or will be) used for diagnostic purposes and to facilitate reflection
to a SYCL run-time library. These properties are not serialized with the AST
but are recreated upon deserialization.

The `sycl_kernel_entry_point` attribute is intended to replace the existing
`sycl_kernel` attribute which is intended to be deprecated in a future change
and removed following an appropriate deprecation period. The new attribute
differs in that it is enabled for both SYCL host and device compilation, may
be used with non-template functions, explicitly indicates the type used as
the kernel name type, and will impact AST generation.

This change adds the basic infrastructure for the new attribute. Future
changes will add diagnostics and new AST support that will be used to drive
generation of the corresponding offload kernel.
---
 clang/include/clang/AST/ASTContext.h          |  12 ++
 clang/include/clang/AST/SYCLKernelInfo.h      |  47 ++++++
 clang/include/clang/Basic/Attr.td             |  16 +-
 clang/include/clang/Basic/AttrDocs.td         |  58 +++++++
 clang/include/clang/Sema/SemaSYCL.h           |   1 +
 clang/lib/AST/ASTContext.cpp                  |  25 +++
 clang/lib/Sema/SemaDecl.cpp                   |   4 +
 clang/lib/Sema/SemaDeclAttr.cpp               |   3 +
 clang/lib/Sema/SemaSYCL.cpp                   |   9 ++
 clang/lib/Serialization/ASTReaderDecl.cpp     |   9 ++
 .../ast-dump-sycl-kernel-entry-point.cpp      | 144 ++++++++++++++++++
 ...a-attribute-supported-attributes-list.test |   1 +
 .../sycl-kernel-entry-point-attr-grammar.cpp  | 137 +++++++++++++++++
 .../sycl-kernel-entry-point-attr-ignored.cpp  |  17 +++
 14 files changed, 480 insertions(+), 3 deletions(-)
 create mode 100644 clang/include/clang/AST/SYCLKernelInfo.h
 create mode 100644 clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
 create mode 100644 clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
 create mode 100644 clang/test/SemaSYCL/sycl-kernel-entry-point-attr-ignored.cpp

diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h
index b65a1f7dff5bc1..e68f0a4da57b16 100644
--- a/clang/include/clang/AST/ASTContext.h
+++ b/clang/include/clang/AST/ASTContext.h
@@ -23,6 +23,7 @@
 #include "clang/AST/ExternalASTSource.h"
 #include "clang/AST/PrettyPrinter.h"
 #include "clang/AST/RawCommentList.h"
+#include "clang/AST/SYCLKernelInfo.h"
 #include "clang/AST/TemplateName.h"
 #include "clang/Basic/LLVM.h"
 #include "clang/Basic/PartialDiagnostic.h"
@@ -1222,6 +1223,11 @@ class ASTContext : public RefCountedBase<ASTContext> {
   /// in device compilation.
   llvm::DenseSet<const FunctionDecl *> CUDAImplicitHostDeviceFunUsedByDevice;
 
+  /// Map of SYCL kernels indexed by the unique type used to name the kernel.
+  /// Entries are not serialized but are recreated on deserialization of a
+  /// sycl_kernel_entry_point attributed function declaration.
+  llvm::DenseMap<CanQualType, SYCLKernelInfo> SYCLKernels;
+
   /// For capturing lambdas with an explicit object parameter whose type is
   /// derived from the lambda type, we need to perform derived-to-base
   /// conversion so we can access the captures; the cast paths for that
@@ -3301,6 +3307,12 @@ class ASTContext : public RefCountedBase<ASTContext> {
   void getFunctionFeatureMap(llvm::StringMap<bool> &FeatureMap,
                              GlobalDecl GD) const;
 
+  /// Generates and stores SYCL kernel metadata for the provided
+  /// SYCL kernel entry point function. The provided function must have
+  /// an attached sycl_kernel_entry_point attribute that specifies a unique
+  /// type for the name of a SYCL kernel.
+  void registerSYCLEntryPointFunction(FunctionDecl *FD);
+
   //===--------------------------------------------------------------------===//
   //                    Statistics
   //===--------------------------------------------------------------------===//
diff --git a/clang/include/clang/AST/SYCLKernelInfo.h b/clang/include/clang/AST/SYCLKernelInfo.h
new file mode 100644
index 00000000000000..79a83330f1d23b
--- /dev/null
+++ b/clang/include/clang/AST/SYCLKernelInfo.h
@@ -0,0 +1,47 @@
+//===--- SYCLKernelInfo.h --- Information about SYCL kernels --------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+/// \file
+/// This file declares types used to describe SYCL kernels.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_CLANG_AST_SYCLKERNELINFO_H
+#define LLVM_CLANG_AST_SYCLKERNELINFO_H
+
+#include <string>
+#include "clang/AST/Decl.h"
+#include "clang/AST/Type.h"
+
+namespace clang {
+
+class SYCLKernelInfo {
+public:
+  SYCLKernelInfo(
+      CanQualType KernelNameType,
+      const FunctionDecl *KernelEntryPointDecl)
+  :
+      KernelNameType(KernelNameType),
+      KernelEntryPointDecl(KernelEntryPointDecl)
+  {}
+
+  CanQualType GetKernelNameType() const {
+    return KernelNameType;
+  }
+
+  const FunctionDecl* GetKernelEntryPointDecl() const {
+    return KernelEntryPointDecl;
+  }
+
+private:
+  CanQualType KernelNameType;
+  const FunctionDecl *KernelEntryPointDecl;
+};
+
+} // namespace clang
+
+#endif // LLVM_CLANG_AST_SYCLKERNELINFO_H
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index ce86116680d7a3..c4a3615752bf10 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -407,7 +407,8 @@ def MicrosoftExt : LangOpt<"MicrosoftExt">;
 def Borland : LangOpt<"Borland">;
 def CUDA : LangOpt<"CUDA">;
 def HIP : LangOpt<"HIP">;
-def SYCL : LangOpt<"SYCLIsDevice">;
+def SYCLHost : LangOpt<"SYCLIsHost">;
+def SYCLDevice : LangOpt<"SYCLIsDevice">;
 def COnly : LangOpt<"", "!LangOpts.CPlusPlus">;
 def CPlusPlus : LangOpt<"CPlusPlus">;
 def OpenCL : LangOpt<"OpenCL">;
@@ -1489,14 +1490,23 @@ def : MutualExclusions<[CUDAConstant, CUDAShared, HIPManaged]>;
 def SYCLKernel : InheritableAttr {
   let Spellings = [Clang<"sycl_kernel">];
   let Subjects = SubjectList<[FunctionTmpl]>;
-  let LangOpts = [SYCL];
+  let LangOpts = [SYCLDevice];
   let Documentation = [SYCLKernelDocs];
 }
 
+def SYCLKernelEntryPoint : InheritableAttr {
+  let Spellings = [Clang<"sycl_kernel_entry_point">];
+  let Args = [TypeArgument<"KernelName">];
+  let Subjects = SubjectList<[Function], ErrorDiag>;
+  let TemplateDependent = 1;
+  let LangOpts = [SYCLHost, SYCLDevice];
+  let Documentation = [SYCLKernelEntryPointDocs];
+}
+
 def SYCLSpecialClass: InheritableAttr {
   let Spellings = [Clang<"sycl_special_class">];
   let Subjects = SubjectList<[CXXRecord]>;
-  let LangOpts = [SYCL];
+  let LangOpts = [SYCLDevice];
   let Documentation = [SYCLSpecialClassDocs];
 }
 
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 8ef151b3f2fddb..cd3aec8f70f024 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -455,6 +455,64 @@ The SYCL kernel in the previous code sample meets these expectations.
   }];
 }
 
+def SYCLKernelEntryPointDocs : Documentation {
+  let Category = DocCatFunction;
+  let Content = [{
+The ``sycl_kernel_entry_point`` attribute specifies that a function definition
+defines a pattern for an offload kernel entry point function to be emitted when
+the source code is compiled with ``-fsycl`` for a device target. Such functions
+serve as the execution entry point for a SYCL run-time library to invoke a SYCL
+kernel on a device. The function's parameters define the parameters to the
+offload kernel.
+
+The attribute requires a single type argument that specifies a class type that
+meets the requirements for a SYCL kernel name as described in section 5.2,
+"Naming of kernels", of the SYCL 2020 specification. A unique kernel name type
+is required for each function declared with the attribute. The attribute may
+not first appear on a declaration that follows a definition of the function.
+
+The attribute appertains only to non-member functions and static member
+functions that meet the following requirements:
+
+- Has a ``void`` return type.
+- Is not a variadic function.
+- Is not a coroutine.
+- Is not defined as deleted or as defaulted.
+- Is not declared with the ``constexpr`` or ``consteval`` specifiers.
+- Is not declared with the ``[[noreturn]]`` attribute.
+
+This attribute is intended for use in the implementation of SYCL run-time
+libraries that implement SYCL kernel invocation functions like the
+``single_task`` and ``parallel_for`` member functions of the ``sycl::handler``
+class specified in section 4.9.4, "Command group ``handler`` class" of the
+SYCL 2020 specification. Such use might look something like the following.
+
+.. code-block:: c++
+
+  namespace sycl {
+  class handler {
+    template<typename KernelNameType, typename KernelType>
+    [[ clang::sycl_kernel_entry_point(KernelNameType) ]]
+    static void kernel_entry_point(KernelType kernel) {
+      kernel();
+    }
+
+  public:
+    template<typename KernelNameType, typename KernelType>
+    void single_task(KernelType kernel) {
+      kernel_entry_point<KernelNameType>(kernel);
+    }
+  };
+  } // namespace sycl
+
+It is not necessary for a SYCL kernel entry point function to be called for
+the offload kernel entry point to be emitted. For inline functions and function
+templates, any ODR-use will suffice. For other functions, an ODR-use is not
+required; the offload kernel entry point will be emitted if the function is
+defined.
+  }];
+}
+
 def SYCLSpecialClassDocs : Documentation {
   let Category = DocCatStmt;
   let Content = [{
diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h
index 27c42b54018307..c9f3358124eda7 100644
--- a/clang/include/clang/Sema/SemaSYCL.h
+++ b/clang/include/clang/Sema/SemaSYCL.h
@@ -62,6 +62,7 @@ class SemaSYCL : public SemaBase {
                                        ParsedType ParsedTy);
 
   void handleKernelAttr(Decl *D, const ParsedAttr &AL);
+  void handleKernelEntryPointAttr(Decl *D, const ParsedAttr &AL);
 };
 
 } // namespace clang
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index 85b3984940ffc2..daf4b8398bdd55 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -14296,6 +14296,31 @@ void ASTContext::getFunctionFeatureMap(llvm::StringMap<bool> &FeatureMap,
   }
 }
 
+static SYCLKernelInfo BuildSYCLKernelInfo(ASTContext &Context,
+                                          CanQualType KernelNameType,
+                                          const FunctionDecl *FD) {
+  return { KernelNameType, FD };
+}
+
+void ASTContext::registerSYCLEntryPointFunction(FunctionDecl *FD) {
+  assert(!FD->isInvalidDecl());
+  assert(!FD->isDependentContext());
+
+  const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>();
+  assert(SKEPAttr && "Missing sycl_kernel_entry_point attribute");
+
+  CanQualType KernelNameType = getCanonicalType(SKEPAttr->getKernelName());
+  auto IT = SYCLKernels.find(KernelNameType);
+  if (IT != SYCLKernels.end()) {
+    if (!declaresSameEntity(FD, IT->second.GetKernelEntryPointDecl()))
+      llvm::report_fatal_error("SYCL kernel name conflict");
+  } else {
+    SYCLKernels.insert_or_assign(
+        KernelNameType,
+        BuildSYCLKernelInfo(*this, KernelNameType, FD));
+  }
+}
+
 OMPTraitInfo &ASTContext::getNewOMPTraitInfo() {
   OMPTraitInfoVector.emplace_back(new OMPTraitInfo());
   return *OMPTraitInfoVector.back();
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index de8805e15bc750..6af5264f79e1f9 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -12053,6 +12053,10 @@ bool Sema::CheckFunctionDeclaration(Scope *S, FunctionDecl *NewFD,
   if (LangOpts.OpenMP)
     OpenMP().ActOnFinishedFunctionDefinitionInOpenMPAssumeScope(NewFD);
 
+  if (LangOpts.isSYCL() && NewFD->hasAttr<SYCLKernelEntryPointAttr>() &&
+      !NewFD->isInvalidDecl() && !NewFD->isDependentContext())
+    getASTContext().registerSYCLEntryPointFunction(NewFD);
+
   // Semantic checking for this function declaration (in isolation).
 
   if (getLangOpts().CPlusPlus) {
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 14cc51cf89665a..2504fc2e7b8ec9 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -6606,6 +6606,9 @@ ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, const ParsedAttr &AL,
   case ParsedAttr::AT_SYCLKernel:
     S.SYCL().handleKernelAttr(D, AL);
     break;
+  case ParsedAttr::AT_SYCLKernelEntryPoint:
+    S.SYCL().handleKernelEntryPointAttr(D, AL);
+    break;
   case ParsedAttr::AT_SYCLSpecialClass:
     handleSimpleAttribute<SYCLSpecialClassAttr>(S, D, AL);
     break;
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index f2d13d456c25fc..a168e654ad6876 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -198,3 +198,12 @@ void SemaSYCL::handleKernelAttr(Decl *D, const ParsedAttr &AL) {
 
   handleSimpleAttribute<SYCLKernelAttr>(*this, D, AL);
 }
+
+void SemaSYCL::handleKernelEntryPointAttr(Decl *D, const ParsedAttr &AL) {
+  ParsedType PT = AL.getTypeArg();
+  TypeSourceInfo *TSI = nullptr;
+  (void)SemaRef.GetTypeFromParser(PT, &TSI);
+  assert(TSI && "no type source info for attribute argument");
+  D->addAttr(::new (SemaRef.Context) SYCLKernelEntryPointAttr(SemaRef.Context,
+                                                              AL, TSI));
+}
diff --git a/clang/lib/Serialization/ASTReaderDecl.cpp b/clang/lib/Serialization/ASTReaderDecl.cpp
index 321e65fd2b094f..02e723aef0ff8a 100644
--- a/clang/lib/Serialization/ASTReaderDecl.cpp
+++ b/clang/lib/Serialization/ASTReaderDecl.cpp
@@ -1155,6 +1155,15 @@ void ASTDeclReader::VisitFunctionDecl(FunctionDecl *FD) {
   for (unsigned I = 0; I != NumParams; ++I)
     Params.push_back(readDeclAs<ParmVarDecl>());
   FD->setParams(Reader.getContext(), Params);
+
+  // If the declaration is a SYCL kernel entry point function as indicated by
+  // the presence of a sycl_kernel_entry_point attribute, register it so that
+  // associated metadata is recreated.
+  if (!FD->isInvalidDecl() && !FD->isDependentContext() &&
+      FD->hasAttr<SYCLKernelEntryPointAttr>()) {
+    ASTContext &C = Reader.getContext();
+    C.registerSYCLEntryPointFunction(FD);
+  }
 }
 
 void ASTDeclReader::VisitObjCMethodDecl(ObjCMethodDecl *MD) {
diff --git a/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp b/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
new file mode 100644
index 00000000000000..342d2f71c357e2
--- /dev/null
+++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
@@ -0,0 +1,144 @@
+// Tests without serialization:
+// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-unknown -fsycl-is-device \
+// RUN:   -ast-dump %s \
+// RUN:   | FileCheck --match-full-lines %s
+// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-unknown -fsycl-is-host \
+// RUN:   -ast-dump %s \
+// RUN:   | FileCheck --match-full-lines %s
+//
+// Tests with serialization:
+// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-unknown -fsycl-is-device \
+// RUN:   -emit-pch -o %t %s
+// RUN: %clang_cc1 -x c++ -std=c++17 -triple x86_64-unknown-unknown -fsycl-is-device \
+// RUN:   -include-pch %t -ast-dump-all /dev/null \
+// RUN:   | sed -e "s/ <undeserialized declarations>//" -e "s/ imported//" \
+// RUN:   | FileCheck --match-full-lines %s
+// RUN: %clang_cc1 -std=c++17 -triple x86_64-unknown-unknown -fsycl-is-host \
+// RUN:   -emit-pch -o %t %s
+// RUN: %clang_cc1 -x c++ -std=c++17 -triple x86_64-unknown-unknown -fsycl-is-host \
+// RUN:   -include-pch %t -ast-dump-all /dev/null \
+// RUN:   | sed -e "s/ <undeserialized declarations>//" -e "s/ imported//" \
+// RUN:   | FileCheck --match-full-lines %s
+
+// These tests validate the AST produced for functions declared with the
+// sycl_kernel_entry_point attribute.
+
+// CHECK: TranslationUnitDecl {{.*}}
+
+// A unique kernel name type is required for each declared kernel entry point.
+template<int, int=0> struct KN;
+
+__attribute__((sycl_kernel_entry_point(KN<1>)))
+void skep1() {
+}
+// CHECK: |-FunctionDecl {{.*}} skep1 'void ()'
+// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} KN<1>
+
+using KN2 = KN<2>;
+__attribute__((sycl_kernel_entry_point(KN2)))
+void skep2() {
+}
+// CHECK: |-FunctionDecl {{.*}} skep2 'void ()'
+// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} KN2
+
+template<int I> using KNT = KN<I>;
+__attribute__((sycl_kernel_entry_point(KNT<3>)))
+void skep3() {
+}
+// CHECK: |-FunctionDecl {{.*}} skep3 'void ()'
+// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} KNT<3>
+
+template<typename KNT, typename F>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void skep4(F f) {
+  f();
+}
+// CHECK: |-FunctionTemplateDecl {{.*}} skep4
+// CHECK: | |-TemplateTypeParmDecl {{.*}} KNT
+// CHECK: | |-TemplateTypeParmDecl {{.*}} F
+// CHECK: | |-FunctionDecl {{.*}} skep4 'void (F)'
+// CHECK: | | `-SYCLKernelEntryPointAttr {{.*}} KNT
+
+void test_skep4() {
+  skep4<KNT<4>>([]{});
+}
+// CHECK: | `-FunctionDecl {{.*}} used skep4 'void ((lambda at {{.*}}))' implicit_instantiation
+// CHECK: |   |-TemplateArgument type 'KN<4>'
+// CHECK: |   |-TemplateArgument type '(lambda at {{.*}})'
+// CHECK: |   `-SYCLKernelEntryPointAttr {{.*}} struct KN<4>
+// CHECK: |-FunctionDecl {{.*}} test_skep4 'void ()'
+
+template<typename KNT, typename T>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void skep5(T) {
+}
+// CHECK: |-FunctionTemplateDecl {{.*}} skep5
+// CHECK: | |-TemplateTypeParmDecl {{.*}} KNT
+// CHECK: | |-TemplateTypeParmDecl {{.*}} T
+// CHECK: | |-FunctionDecl {{.*}} skep5 'void (T)'
+// CHECK: | | `-SYCLKernelEntryPointAttr {{.*}} KNT
+
+// Checks for the explicit template instantiation declaration below.
+// CHECK: | `-FunctionDecl {{.*}} skep5 'void (int)' explicit_instantiation_definition
+// CHECK: |   |-TemplateArgument type 'KN<5, 4>'
+// CHECK: |   |-TemplateArgument type 'int'
+// CHECK: |   `-SYCLKernelEntryPointAttr {{.*}} KN<5, 4>
+
+// FIXME: C++23 [temp.expl.spec]p12 states:
+// FIXME:   ... Similarly, attributes appearing in the declaration of a template
+// FIXME:   have no effect on an explicit specialization of that template.
+// FIXME: Clang currently instantiates and propagates attributes from a function
+// FIXME: template to its explicit specializations resulting in the following
+// FIXME: explicit specialization having an attribute incorrectly attached.
+template<>
+void skep5<KN<5,1>>(short) {
+}
+// CHECK: |-FunctionDecl {{.*}} prev {{.*}} skep5 'void (short)' explicit_specialization
+// CHECK: | |-TemplateArgument type 'KN<5, 1>'
+// CHECK: | |-TemplateArgument type 'short'
+// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} Inherited struct KN<5, 1>
+
+template<>
+[[clang::sycl_kernel_entry_point(KN<5,2>)]]
+void skep5<KN<5,2>>(long) {
+}
+// CHECK: |-FunctionDecl {{.*}} prev {{.*}} skep5 'void (long)' explicit_specialization
+// CHECK: | |-TemplateArgument type 'KN<5, 2>'
+// CHECK: | |-TemplateArgument type 'long'
+// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} KN<5, 2>
+
+template<>
+[[clang::sycl_kernel_entry_point(KN<5,3>)]]
+void skep5<KN<5,-1>>(long long) {
+}
+// CHECK: |-FunctionDecl {{.*}} prev {{.*}} skep5 'void (long long)' explicit_specialization
+// CHECK: | |-TemplateArgument type 'KN<5, -1>'
+// CHECK: | |-TemplateArgument type 'long long'
+// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} KN<5, 3>
+
+template void skep5<KN<5,4>>(int);
+// Checks are located with the primary template declaration above.
+
+// Ensure that matching attributes from multiple declarations are ok.
+[[clang::sycl_kernel_entry_point(KN<6>)]]
+void skep6();
+[[clang::sycl_kernel_entry_point(KN<6>)]]
+void skep6() {
+}
+// CHECK:      |-FunctionDecl {{.*}} skep6 'void ()'
+// CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} KN<6>
+// CHECK:      |-FunctionDecl {{.*}} prev {{.*}} skep6 'void ()'
+// CHECK:      | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} KN<6>
+
+// Ensure that matching attributes from the same declaration are ok.
+[[clang::sycl_kernel_entry_point(KN<7>), clang::sycl_kernel_entry_point(KN<7>)]]
+void skep7() {
+}
+// CHECK:      |-FunctionDecl {{.*}} skep7 'void ()'
+// CHECK:      | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | |-SYCLKernelEntryPointAttr {{.*}} KN<7>
+// CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} KN<7>
+
+void the_end() {}
+// CHECK: `-FunctionDecl {{.*}} the_end 'void ()'
diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
index 914f94c08a9fd9..5c2f3a347dfb79 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -179,6 +179,7 @@
 // CHECK-NEXT: ReturnTypestate (SubjectMatchRule_function, SubjectMatchRule_variable_is_parameter)
 // CHECK-NEXT: ReturnsNonNull (SubjectMatchRule_objc_method, SubjectMatchRule_function)
 // CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function)
+// CHECK-NEXT: SYCLKernelEntryPoint (SubjectMatchRule_function)
 // CHECK-NEXT: SYCLSpecialClass (SubjectMatchRule_record)
 // CHECK-NEXT: ScopedLockable (SubjectMatchRule_record)
 // CHECK-NEXT: Section (SubjectMatchRule_function, SubjectMatchRule_variable_is_global, SubjectMatchRule_objc_method, SubjectMatchRule_objc_property)
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
new file mode 100644
index 00000000000000..c63d241163e618
--- /dev/null
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-grammar.cpp
@@ -0,0 +1,137 @@
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -fsycl-is-device -verify %s
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++20 -fsyntax-only -fsycl-is-device -verify %s
+
+// These tests validate parsing of the sycl_kernel_entry_point argument list
+// and that the single argument names a type.
+
+// Templates used to exercise class template specializations.
+template<int> struct ST; // #ST-decl
+template<int N> using TTA = ST<N>; // #TTA-decl
+
+
+////////////////////////////////////////////////////////////////////////////////
+// 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() {}
+
+template<typename T>
+[[clang::sycl_kernel_entry_point(T)]] void ok13(T k);
+void test_ok13() {
+  ok13([]{});
+}
+
+
+////////////////////////////////////////////////////////////////////////////////
+// Invalid declarations.
+////////////////////////////////////////////////////////////////////////////////
+
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute takes one argument}}
+[[clang::sycl_kernel_entry_point]] void bad1();
+
+// expected-error@+1 {{'sycl_kernel_entry_point' attribute takes one argument}}
+[[clang::sycl_kernel_entry_point()]] void bad2();
+
+struct B3;
+// expected-error@+2 {{expected ')'}}
+// expected-error@+1 {{expected ']'}}
+[[clang::sycl_kernel_entry_point(B3,)]] void bad3();
+
+struct B4;
+// expected-error@+3 {{expected ')'}}
+// expected-error@+2 {{expected ','}}
+// expected-warning@+1 {{unknown attribute 'X' ignored}}
+[[clang::sycl_kernel_entry_point(B4, X)]] void bad4();
+
+// expected-error@+1 {{expected a type}}
+[[clang::sycl_kernel_entry_point(1)]] void bad5();
+
+void f6();
+// expected-error@+1 {{unknown type name 'f6'}}
+[[clang::sycl_kernel_entry_point(f6)]] void bad6();
+
+// expected-error@+2 {{use of class template 'ST' requires template arguments; argument deduction not allowed here}}
+// expected-note@#ST-decl {{template is declared here}}
+[[clang::sycl_kernel_entry_point(ST)]] void bad7();
+
+// expected-error@+2 {{use of alias template 'TTA' requires template arguments; argument deduction not allowed here}}
+// expected-note@#TTA-decl {{template is declared here}}
+[[clang::sycl_kernel_entry_point(TTA)]] void bad8();
+
+enum {
+  e9
+};
+// expected-error@+1 {{unknown type name 'e9'}}
+[[clang::sycl_kernel_entry_point(e9)]] void bad9();
+
+#if __cplusplus >= 202002L
+template<typename> concept C = true;
+// expected-error@+1 {{expected a type}}
+[[clang::sycl_kernel_entry_point(C)]] void bad10();
+
+// expected-error@+1 {{expected a type}}
+[[clang::sycl_kernel_entry_point(C<int>)]] void bad11();
+#endif
+
+struct B12; // #B12-decl
+// FIXME: C++23 [temp.expl.spec]p12 states:
+// FIXME:   ... Similarly, attributes appearing in the declaration of a template
+// FIXME:   have no effect on an explicit specialization of that template.
+// FIXME: Clang currently instantiates and propagates attributes from a function
+// FIXME: template to its explicit specializations resulting in the following
+// FIXME: spurious error.
+// expected-error@+4 {{incomplete type 'B12' named in nested name specifier}}
+// expected-note@+5 {{in instantiation of function template specialization 'bad12<B12>' requested here}}
+// expected-note@#B12-decl {{forward declaration of 'B12'}}
+template<typename T>
+[[clang::sycl_kernel_entry_point(typename T::not_found)]] void bad12() {}
+template<>
+void bad12<B12>() {}
diff --git a/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-ignored.cpp b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-ignored.cpp
new file mode 100644
index 00000000000000..30de6ae0b0e6f9
--- /dev/null
+++ b/clang/test/SemaSYCL/sycl-kernel-entry-point-attr-ignored.cpp
@@ -0,0 +1,17 @@
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -verify %s
+
+// These tests validate that the sycl_kernel_entry_point attribute is ignored
+// when SYCL support is not enabled.
+
+// A unique kernel name type is required for each declared kernel entry point.
+template<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>>();

From 870a49e29fab88ac033cb94bdfd3dc87c1c50685 Mon Sep 17 00:00:00 2001
From: Tom Honermann <tom.honermann@intel.com>
Date: Mon, 7 Oct 2024 14:25:31 -0700
Subject: [PATCH 2/9] Fix clang-format complaints.

---
 clang/include/clang/AST/SYCLKernelInfo.h | 21 ++++++++-------------
 clang/lib/AST/ASTContext.cpp             |  5 ++---
 clang/lib/Sema/SemaSYCL.cpp              |  4 ++--
 3 files changed, 12 insertions(+), 18 deletions(-)

diff --git a/clang/include/clang/AST/SYCLKernelInfo.h b/clang/include/clang/AST/SYCLKernelInfo.h
index 79a83330f1d23b..7a88d5ec70dca6 100644
--- a/clang/include/clang/AST/SYCLKernelInfo.h
+++ b/clang/include/clang/AST/SYCLKernelInfo.h
@@ -13,27 +13,22 @@
 #ifndef LLVM_CLANG_AST_SYCLKERNELINFO_H
 #define LLVM_CLANG_AST_SYCLKERNELINFO_H
 
-#include <string>
 #include "clang/AST/Decl.h"
 #include "clang/AST/Type.h"
+#include <string>
 
 namespace clang {
 
 class SYCLKernelInfo {
 public:
-  SYCLKernelInfo(
-      CanQualType KernelNameType,
-      const FunctionDecl *KernelEntryPointDecl)
-  :
-      KernelNameType(KernelNameType),
-      KernelEntryPointDecl(KernelEntryPointDecl)
-  {}
-
-  CanQualType GetKernelNameType() const {
-    return KernelNameType;
-  }
+  SYCLKernelInfo(CanQualType KernelNameType,
+                 const FunctionDecl *KernelEntryPointDecl)
+      : KernelNameType(KernelNameType),
+        KernelEntryPointDecl(KernelEntryPointDecl) {}
+
+  CanQualType GetKernelNameType() const { return KernelNameType; }
 
-  const FunctionDecl* GetKernelEntryPointDecl() const {
+  const FunctionDecl *GetKernelEntryPointDecl() const {
     return KernelEntryPointDecl;
   }
 
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index daf4b8398bdd55..397cb9330923d1 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -14299,7 +14299,7 @@ void ASTContext::getFunctionFeatureMap(llvm::StringMap<bool> &FeatureMap,
 static SYCLKernelInfo BuildSYCLKernelInfo(ASTContext &Context,
                                           CanQualType KernelNameType,
                                           const FunctionDecl *FD) {
-  return { KernelNameType, FD };
+  return {KernelNameType, FD};
 }
 
 void ASTContext::registerSYCLEntryPointFunction(FunctionDecl *FD) {
@@ -14316,8 +14316,7 @@ void ASTContext::registerSYCLEntryPointFunction(FunctionDecl *FD) {
       llvm::report_fatal_error("SYCL kernel name conflict");
   } else {
     SYCLKernels.insert_or_assign(
-        KernelNameType,
-        BuildSYCLKernelInfo(*this, KernelNameType, FD));
+        KernelNameType, BuildSYCLKernelInfo(*this, KernelNameType, FD));
   }
 }
 
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index a168e654ad6876..e7cecebae25808 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -204,6 +204,6 @@ void SemaSYCL::handleKernelEntryPointAttr(Decl *D, const ParsedAttr &AL) {
   TypeSourceInfo *TSI = nullptr;
   (void)SemaRef.GetTypeFromParser(PT, &TSI);
   assert(TSI && "no type source info for attribute argument");
-  D->addAttr(::new (SemaRef.Context) SYCLKernelEntryPointAttr(SemaRef.Context,
-                                                              AL, TSI));
+  D->addAttr(::new (SemaRef.Context)
+                 SYCLKernelEntryPointAttr(SemaRef.Context, AL, TSI));
 }

From 0b57d81189b7bd49c1ced07875ad75b59ba6b4eb Mon Sep 17 00:00:00 2001
From: Tom Honermann <tom.honermann@intel.com>
Date: Fri, 11 Oct 2024 18:25:15 -0700
Subject: [PATCH 3/9] Address code review feedback from Erich Keane.

---
 clang/include/clang/AST/SYCLKernelInfo.h |   4 +-
 clang/include/clang/Basic/AttrDocs.td    | 158 +++++++++++++++++++----
 clang/lib/AST/ASTContext.cpp             |   6 +-
 3 files changed, 139 insertions(+), 29 deletions(-)

diff --git a/clang/include/clang/AST/SYCLKernelInfo.h b/clang/include/clang/AST/SYCLKernelInfo.h
index 7a88d5ec70dca6..84a5e5e4926d3d 100644
--- a/clang/include/clang/AST/SYCLKernelInfo.h
+++ b/clang/include/clang/AST/SYCLKernelInfo.h
@@ -26,9 +26,9 @@ class SYCLKernelInfo {
       : KernelNameType(KernelNameType),
         KernelEntryPointDecl(KernelEntryPointDecl) {}
 
-  CanQualType GetKernelNameType() const { return KernelNameType; }
+  CanQualType getKernelNameType() const { return KernelNameType; }
 
-  const FunctionDecl *GetKernelEntryPointDecl() const {
+  const FunctionDecl *getKernelEntryPointDecl() const {
     return KernelEntryPointDecl;
   }
 
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index cd3aec8f70f024..e651ab2bb71eba 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -458,12 +458,13 @@ 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 ``sycl_kernel_entry_point`` attribute facilitates the generation of an
+offload kernel entry point, sometimes called a SYCL kernel caller function,
+suitable for invoking a SYCL kernel on an offload device. The attribute is
+intended for use in the implementation of SYCL kernel invocation functions
+like the ``single_task`` and ``parallel_for`` member functions of the
+``sycl::handler`` class specified in section 4.9.4, "Command group ``handler``
+class", of the SYCL 2020 specification.
 
 The attribute requires a single type argument that specifies a class type that
 meets the requirements for a SYCL kernel name as described in section 5.2,
@@ -471,21 +472,19 @@ meets the requirements for a SYCL kernel name as described in section 5.2,
 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:
+The attribute only appertains to functions and only those 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.
+* Has a ``void`` return type.
+* Is not a non-static member function, constructor, or destructor.
+* Is not a C variadic function.
+* Is not a coroutine.
+* Is not defined as deleted or as defaulted.
+* Is not declared with the ``constexpr`` or ``consteval`` specifiers.
+* Is not declared with the ``[[noreturn]]`` attribute.
 
-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.
+Use in the implementation of a SYCL kernel invocation function might look as
+follows.
 
 .. code-block:: c++
 
@@ -500,16 +499,127 @@ SYCL 2020 specification. Such use might look something like the following.
   public:
     template<typename KernelNameType, typename KernelType>
     void single_task(KernelType kernel) {
+      // Call kernel_entry_point() to trigger generation of an offload
+      // kernel entry point.
       kernel_entry_point<KernelNameType>(kernel);
+      // Call functions appropriate for the desired offload backend
+      // (OpenCL, CUDA, HIP, Level Zero, etc...).
     }
   };
   } // 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.
+A SYCL kernel is a callable object of class type that is constructed on a host,
+often via a lambda expression, and then passed to a SYCL kernel invocation
+function to be executed on an offload device. A SYCL kernel invocation function
+is responsible for copying the provided SYCL kernel object to an offload
+device and initiating a call to it. The SYCL kernel object and its data members
+constitute the parameters of an offload kernel.
+
+A SYCL kernel type is required to satisfy the device copyability requirements
+specified in section 3.13.1, "Device copyable", of the SYCL 2020 specification.
+Additionally, any data members of the kernel object type are required to satisfy
+section 4.12.4, "Rules for parameter passing to kernels". For most types, these
+rules require that the type is trivially copyable.  However, the SYCL
+specification mandates that certain special SYCL types, such as
+``sycl::accessor`` and ``sycl::stream`` be device copyable even if they are not
+trivially copyable. These types require special handling because they cannot
+be copied to device memory as if by ``memcpy()``. Additionally, some offload
+backends, OpenCL for example, require objects of some of these types to be
+passed as individual arguments to the offload kernel.
+
+An offload kernel consists of an entry point function that declares the
+parameters of the offload kernel and the set of all functions and variables that
+are directly or indirectly used by the entry point function.
+
+A SYCL kernel invocation function invokes a SYCL kernel on a device by
+performing the following tasks (likely with the help of an offload backend
+like OpenCL):
+
+#. Identifying the offload kernel entry point to be used for the SYCL kernel.
+
+#. Deconstructing the SYCL kernel object, if necessary, to produce the set of
+   offload kernel arguments required by the offload kernel entry point.
+
+#. Copying the offload kernel arguments to device memory.
+
+#. Initiating execution of the offload kernel entry point.
+
+The offload kernel entry point for a SYCL kernel performs the following tasks:
+
+#. Reconstituting the SYCL kernel object, if necessary, using the offload
+   kernel parameters.
+
+#. Calling the ``operator()`` member function of the (reconstituted) SYCL kernel
+   object.
+
+The ``sycl_kernel_entry_point`` attribute automates generation of an offload
+kernel entry point that performs those latter tasks. The parameters and body of
+a function declared with the ``sycl_kernel_entry_point`` attribute specify a
+pattern from which the parameters and body of the entry point function are
+derived. Consider the following call to a SYCL kernel invocation function.
+
+.. code-block:: c++
+
+  struct S { int i; };
+  void f(sycl::handler &handler, sycl::stream &sout, S s) {
+    handler.single_task<struct KN>([=] {
+      sout << "The value of s.i is " << s.i << "\n";
+    });
+  }
+
+The SYCL kernel object is the result of the lambda expression. It has two
+data members corresponding to the captures of ``sout`` and ``s``. Since one
+of these data members corresponds to a special SYCL type that must be passed
+individually as an offload kernel parameter, it is necessary to decompose the
+SYCL kernel object into its constituent parts; the offload kernel will have
+two kernel parameters. Given a SYCL implementation that uses a
+``sycl_kernel_entry_point`` attributed function like the one shown above, an
+offload kernel entry point function will be generated that looks approximately
+as follows.
+
+.. code-block:: c++
+
+  void sycl-kernel-caller-for-KN(sycl::stream sout, S s) {
+    kernel-type kernel = { sout, s );
+    kernel();
+  }
+
+There are a few items worthy of note:
+
+#. The name of the generated function incorporates the SYCL kernel name,
+   ``KN``, that was passed as the ``KernelNameType`` template parameter to
+   ``kernel_entry_point()`` and provided as the argument to the
+   ``sycl_kernel_entry_point`` attribute. There is a one-to-one correspondence
+   between SYCL kernel names and offload kernel entry points.
+
+#. The SYCL kernel is a lambda closure type and therefore has no name;
+   ``kernel-type`` is substituted above and corresponds to the ``KernelType``
+   template parameter deduced in the call to ``kernel_entry_point()``.
+   Lambda types cannot be declared and initialized using the aggregate
+   initialization syntax used above, but the intended behavior should be clear.
+
+#. ``S`` is a device copyable type that does not directly or indirectly contain
+   a data member of a SYCL special type. It therefore does not need to be
+   decomposed into its constituent members to be passed as a kernel argument.
+
+#. The call to ``kernel_entry_point()`` has no effect other than to trigger
+   emission of the entry point function. The statments that make up the body
+   of the function are not executed when the function is called; they are
+   only used in the generation of the entry point function.
+
+It is not necessary for a function declared with the ``sycl_kernel_entry_point``
+attribute to be called for the offload kernel entry point to be emitted. For
+inline functions and function templates, any ODR-use will suffice. For other
+functions, an ODR-use is not required; the offload kernel entry point will be
+emitted if the function is defined.
+
+Functions declared with the ``sycl_kernel_entry_point`` attribute are not
+limited to the simple example shown above. They may have additional template
+parameters, declare additional function parameters, and have complex control
+flow in the function body. Function parameter decomposition and reconstitution
+is performed for all function parameters. The function must abide by the
+language feature restrictions described in section 5.4, "Language restrictions
+for device functions" in the SYCL 2020 specification.
   }];
 }
 
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index 397cb9330923d1..c40e4c384229a1 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -14312,11 +14312,11 @@ void ASTContext::registerSYCLEntryPointFunction(FunctionDecl *FD) {
   CanQualType KernelNameType = getCanonicalType(SKEPAttr->getKernelName());
   auto IT = SYCLKernels.find(KernelNameType);
   if (IT != SYCLKernels.end()) {
-    if (!declaresSameEntity(FD, IT->second.GetKernelEntryPointDecl()))
+    if (!declaresSameEntity(FD, IT->second.getKernelEntryPointDecl()))
       llvm::report_fatal_error("SYCL kernel name conflict");
   } else {
-    SYCLKernels.insert_or_assign(
-        KernelNameType, BuildSYCLKernelInfo(*this, KernelNameType, FD));
+    SYCLKernels.insert(std::make_pair(
+        KernelNameType, BuildSYCLKernelInfo(*this, KernelNameType, FD)));
   }
 }
 

From 0ad8348c2382d71c514c9bca7c69717167f10753 Mon Sep 17 00:00:00 2001
From: Tom Honermann <tom.honermann@intel.com>
Date: Sun, 27 Oct 2024 09:20:10 -0700
Subject: [PATCH 4/9] Address code review feedback from Alexey Bader.

---
 clang/lib/AST/ASTContext.cpp | 5 ++---
 1 file changed, 2 insertions(+), 3 deletions(-)

diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index c40e4c384229a1..84c03bd59f7fac 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -14296,8 +14296,7 @@ void ASTContext::getFunctionFeatureMap(llvm::StringMap<bool> &FeatureMap,
   }
 }
 
-static SYCLKernelInfo BuildSYCLKernelInfo(ASTContext &Context,
-                                          CanQualType KernelNameType,
+static SYCLKernelInfo BuildSYCLKernelInfo(CanQualType KernelNameType,
                                           const FunctionDecl *FD) {
   return {KernelNameType, FD};
 }
@@ -14316,7 +14315,7 @@ void ASTContext::registerSYCLEntryPointFunction(FunctionDecl *FD) {
       llvm::report_fatal_error("SYCL kernel name conflict");
   } else {
     SYCLKernels.insert(std::make_pair(
-        KernelNameType, BuildSYCLKernelInfo(*this, KernelNameType, FD)));
+        KernelNameType, BuildSYCLKernelInfo(KernelNameType, FD)));
   }
 }
 

From c2f24ed3ee581063026893de564b3ee65b76e8ce Mon Sep 17 00:00:00 2001
From: Tom Honermann <tom.honermann@intel.com>
Date: Sun, 27 Oct 2024 10:42:33 -0700
Subject: [PATCH 5/9] Fix clang-format complaints.

---
 clang/lib/AST/ASTContext.cpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index 84c03bd59f7fac..778c62daee1d16 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -14314,8 +14314,8 @@ void ASTContext::registerSYCLEntryPointFunction(FunctionDecl *FD) {
     if (!declaresSameEntity(FD, IT->second.getKernelEntryPointDecl()))
       llvm::report_fatal_error("SYCL kernel name conflict");
   } else {
-    SYCLKernels.insert(std::make_pair(
-        KernelNameType, BuildSYCLKernelInfo(KernelNameType, FD)));
+    SYCLKernels.insert(std::make_pair(KernelNameType,
+                                      BuildSYCLKernelInfo(KernelNameType, FD)));
   }
 }
 

From b9c910924fbecefbb26cc008bba73f6d1de2d6c0 Mon Sep 17 00:00:00 2001
From: Tom Honermann <tom.honermann@intel.com>
Date: Wed, 30 Oct 2024 14:41:50 -0700
Subject: [PATCH 6/9] Round two of addressing code review feedback from Erich
 Keane.

Changes include:
- Removal of an unneeded inclusion of <string>.
- An update to the `sycl_kernel_entry_point` attribute documentation.
- Updates to the AST test to validate next line proximity where applicable.
---
 clang/include/clang/AST/SYCLKernelInfo.h      |  1 -
 clang/include/clang/Basic/AttrDocs.td         |  6 ++
 .../ast-dump-sycl-kernel-entry-point.cpp      | 82 +++++++++----------
 3 files changed, 47 insertions(+), 42 deletions(-)

diff --git a/clang/include/clang/AST/SYCLKernelInfo.h b/clang/include/clang/AST/SYCLKernelInfo.h
index 84a5e5e4926d3d..55dba1f8e31fd9 100644
--- a/clang/include/clang/AST/SYCLKernelInfo.h
+++ b/clang/include/clang/AST/SYCLKernelInfo.h
@@ -15,7 +15,6 @@
 
 #include "clang/AST/Decl.h"
 #include "clang/AST/Type.h"
-#include <string>
 
 namespace clang {
 
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index e651ab2bb71eba..87782146ae14ba 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -602,6 +602,12 @@ There are a few items worthy of note:
    a data member of a SYCL special type. It therefore does not need to be
    decomposed into its constituent members to be passed as a kernel argument.
 
+#. The depiction of the ``sycl::stream`` parameter as a single self contained
+   kernel parameter is an oversimplification. SYCL special types may require
+   additional decomposition such that the generated function might have three
+   or more parameters depending on how the SYCL library implementation defines
+   these types.
+
 #. The call to ``kernel_entry_point()`` has no effect other than to trigger
    emission of the entry point function. The statments that make up the body
    of the function are not executed when the function is called; they are
diff --git a/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp b/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
index 342d2f71c357e2..c351f3b7d03eab 100644
--- a/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
+++ b/clang/test/ASTSYCL/ast-dump-sycl-kernel-entry-point.cpp
@@ -31,58 +31,58 @@ template<int, int=0> struct KN;
 __attribute__((sycl_kernel_entry_point(KN<1>)))
 void skep1() {
 }
-// CHECK: |-FunctionDecl {{.*}} skep1 'void ()'
-// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} KN<1>
+// 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
+// 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>
+// 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
+// CHECK:      |-FunctionTemplateDecl {{.*}} skep4
+// CHECK-NEXT: | |-TemplateTypeParmDecl {{.*}} KNT
+// CHECK-NEXT: | |-TemplateTypeParmDecl {{.*}} F
+// CHECK-NEXT: | |-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 ()'
+// CHECK:      | `-FunctionDecl {{.*}} used skep4 'void ((lambda at {{.*}}))' implicit_instantiation
+// CHECK-NEXT: |   |-TemplateArgument type 'KN<4>'
+// CHECK:      |   |-TemplateArgument type '(lambda at {{.*}})'
+// CHECK:      |   `-SYCLKernelEntryPointAttr {{.*}} struct KN<4>
+// CHECK-NEXT: |-FunctionDecl {{.*}} test_skep4 'void ()'
 
 template<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
+// CHECK:      |-FunctionTemplateDecl {{.*}} skep5
+// CHECK-NEXT: | |-TemplateTypeParmDecl {{.*}} KNT
+// CHECK-NEXT: | |-TemplateTypeParmDecl {{.*}} T
+// CHECK-NEXT: | |-FunctionDecl {{.*}} skep5 'void (T)'
+// CHECK:      | | `-SYCLKernelEntryPointAttr {{.*}} KNT
 
 // Checks for the explicit template instantiation declaration below.
-// CHECK: | `-FunctionDecl {{.*}} skep5 'void (int)' explicit_instantiation_definition
-// CHECK: |   |-TemplateArgument type 'KN<5, 4>'
-// CHECK: |   |-TemplateArgument type 'int'
-// CHECK: |   `-SYCLKernelEntryPointAttr {{.*}} KN<5, 4>
+// CHECK:      | `-FunctionDecl {{.*}} skep5 'void (int)' explicit_instantiation_definition
+// CHECK-NEXT: |   |-TemplateArgument type 'KN<5, 4>'
+// CHECK:      |   |-TemplateArgument type 'int'
+// CHECK:      |   `-SYCLKernelEntryPointAttr {{.*}} KN<5, 4>
 
 // FIXME: C++23 [temp.expl.spec]p12 states:
 // FIXME:   ... Similarly, attributes appearing in the declaration of a template
@@ -93,28 +93,28 @@ void skep5(T) {
 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>
+// CHECK:      |-FunctionDecl {{.*}} prev {{.*}} skep5 'void (short)' explicit_specialization
+// CHECK-NEXT: | |-TemplateArgument type 'KN<5, 1>'
+// CHECK:      | |-TemplateArgument type 'short'
+// CHECK:      | `-SYCLKernelEntryPointAttr {{.*}} Inherited struct KN<5, 1>
 
 template<>
 [[clang::sycl_kernel_entry_point(KN<5,2>)]]
 void skep5<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>
+// CHECK:      |-FunctionDecl {{.*}} prev {{.*}} skep5 'void (long)' explicit_specialization
+// CHECK-NEXT: | |-TemplateArgument type 'KN<5, 2>'
+// CHECK:      | |-TemplateArgument type 'long'
+// CHECK:      | `-SYCLKernelEntryPointAttr {{.*}} KN<5, 2>
 
 template<>
 [[clang::sycl_kernel_entry_point(KN<5,3>)]]
 void skep5<KN<5,-1>>(long long) {
 }
-// CHECK: |-FunctionDecl {{.*}} prev {{.*}} skep5 'void (long long)' explicit_specialization
-// CHECK: | |-TemplateArgument type 'KN<5, -1>'
-// CHECK: | |-TemplateArgument type 'long long'
-// CHECK: | `-SYCLKernelEntryPointAttr {{.*}} KN<5, 3>
+// CHECK:      |-FunctionDecl {{.*}} prev {{.*}} skep5 'void (long long)' explicit_specialization
+// CHECK-NEXT: | |-TemplateArgument type 'KN<5, -1>'
+// CHECK:      | |-TemplateArgument type 'long long'
+// CHECK:      | `-SYCLKernelEntryPointAttr {{.*}} KN<5, 3>
 
 template void skep5<KN<5,4>>(int);
 // Checks are located with the primary template declaration above.
@@ -127,8 +127,8 @@ void skep6() {
 }
 // CHECK:      |-FunctionDecl {{.*}} skep6 'void ()'
 // CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} KN<6>
-// CHECK:      |-FunctionDecl {{.*}} prev {{.*}} skep6 'void ()'
-// CHECK:      | |-CompoundStmt {{.*}}
+// CHECK-NEXT: |-FunctionDecl {{.*}} prev {{.*}} skep6 'void ()'
+// CHECK-NEXT: | |-CompoundStmt {{.*}}
 // CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} KN<6>
 
 // Ensure that matching attributes from the same declaration are ok.
@@ -136,9 +136,9 @@ void skep6() {
 void skep7() {
 }
 // CHECK:      |-FunctionDecl {{.*}} skep7 'void ()'
-// CHECK:      | |-CompoundStmt {{.*}}
+// CHECK-NEXT: | |-CompoundStmt {{.*}}
 // CHECK-NEXT: | |-SYCLKernelEntryPointAttr {{.*}} KN<7>
 // CHECK-NEXT: | `-SYCLKernelEntryPointAttr {{.*}} KN<7>
 
 void the_end() {}
-// CHECK: `-FunctionDecl {{.*}} the_end 'void ()'
+// CHECK:      `-FunctionDecl {{.*}} the_end 'void ()'

From 732309b501167ea7f48a2dcb03b9bcf0d2996a88 Mon Sep 17 00:00:00 2001
From: Tom Honermann <tom.honermann@intel.com>
Date: Fri, 1 Nov 2024 10:03:06 -0700
Subject: [PATCH 7/9] Round three of addressing code review feedback from Erich
 Keane.

Changes include:
- Use `isTemplated()` instead of `isDependentContext()` since the latter doesn't
  necessarily handle templated friend functions.
---
 clang/lib/AST/ASTContext.cpp              | 2 +-
 clang/lib/Sema/SemaDecl.cpp               | 2 +-
 clang/lib/Serialization/ASTReaderDecl.cpp | 2 +-
 3 files changed, 3 insertions(+), 3 deletions(-)

diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index 778c62daee1d16..03648b355a255e 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -14303,7 +14303,7 @@ static SYCLKernelInfo BuildSYCLKernelInfo(CanQualType KernelNameType,
 
 void ASTContext::registerSYCLEntryPointFunction(FunctionDecl *FD) {
   assert(!FD->isInvalidDecl());
-  assert(!FD->isDependentContext());
+  assert(!FD->isTemplated());
 
   const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>();
   assert(SKEPAttr && "Missing sycl_kernel_entry_point attribute");
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 6af5264f79e1f9..cc5ef7b1502e79 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -12054,7 +12054,7 @@ bool Sema::CheckFunctionDeclaration(Scope *S, FunctionDecl *NewFD,
     OpenMP().ActOnFinishedFunctionDefinitionInOpenMPAssumeScope(NewFD);
 
   if (LangOpts.isSYCL() && NewFD->hasAttr<SYCLKernelEntryPointAttr>() &&
-      !NewFD->isInvalidDecl() && !NewFD->isDependentContext())
+      !NewFD->isInvalidDecl() && !NewFD->isTemplated())
     getASTContext().registerSYCLEntryPointFunction(NewFD);
 
   // Semantic checking for this function declaration (in isolation).
diff --git a/clang/lib/Serialization/ASTReaderDecl.cpp b/clang/lib/Serialization/ASTReaderDecl.cpp
index 02e723aef0ff8a..7e5c08b41d5847 100644
--- a/clang/lib/Serialization/ASTReaderDecl.cpp
+++ b/clang/lib/Serialization/ASTReaderDecl.cpp
@@ -1159,7 +1159,7 @@ void ASTDeclReader::VisitFunctionDecl(FunctionDecl *FD) {
   // 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() &&
+  if (!FD->isInvalidDecl() && !FD->isTemplated() &&
       FD->hasAttr<SYCLKernelEntryPointAttr>()) {
     ASTContext &C = Reader.getContext();
     C.registerSYCLEntryPointFunction(FD);

From 8a51f1cde97e2125e2d7bbf84f6631f96ab88b90 Mon Sep 17 00:00:00 2001
From: Tom Honermann <tom.honermann@intel.com>
Date: Fri, 1 Nov 2024 15:32:39 -0700
Subject: [PATCH 8/9] Round four of addressing code review feedback from Erich
 Keane.

Changes include:
- Replaced a use of `llvm::report_fatal_error()` with `assert()`.
- Modified `registerSYCLEntryPointFunction()` to silently skip registration
  of invalid or dependent function declarations.
---
 clang/include/clang/AST/ASTContext.h      |  4 +++-
 clang/lib/AST/ASTContext.cpp              | 13 +++++++++----
 clang/lib/Sema/SemaDecl.cpp               |  3 +--
 clang/lib/Serialization/ASTReaderDecl.cpp |  3 +--
 4 files changed, 14 insertions(+), 9 deletions(-)

diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h
index e68f0a4da57b16..eae80aed479f8d 100644
--- a/clang/include/clang/AST/ASTContext.h
+++ b/clang/include/clang/AST/ASTContext.h
@@ -3310,7 +3310,9 @@ class ASTContext : public RefCountedBase<ASTContext> {
   /// 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.
+  /// type for the name of a SYCL kernel. Callers are required to detect
+  /// conflicting SYCL kernel names and issue a diagnostic prior to calling
+  /// this function.
   void registerSYCLEntryPointFunction(FunctionDecl *FD);
 
   //===--------------------------------------------------------------------===//
diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index 03648b355a255e..b7f933e6e2f746 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -14302,17 +14302,22 @@ static SYCLKernelInfo BuildSYCLKernelInfo(CanQualType KernelNameType,
 }
 
 void ASTContext::registerSYCLEntryPointFunction(FunctionDecl *FD) {
-  assert(!FD->isInvalidDecl());
-  assert(!FD->isTemplated());
+  // If the function declaration to register is invalid or dependent, the
+  // registration attempt is ignored.
+  if (FD->isInvalidDecl() || FD->isTemplated())
+    return;
 
   const auto *SKEPAttr = FD->getAttr<SYCLKernelEntryPointAttr>();
   assert(SKEPAttr && "Missing sycl_kernel_entry_point attribute");
 
+  // Be tolerant of multiple registration attempts so long as each attempt
+  // is for the same entity. Callers are obligated to detect and diagnose
+  // conflicting kernel names prior to calling this function.
   CanQualType KernelNameType = getCanonicalType(SKEPAttr->getKernelName());
   auto IT = SYCLKernels.find(KernelNameType);
   if (IT != SYCLKernels.end()) {
-    if (!declaresSameEntity(FD, IT->second.getKernelEntryPointDecl()))
-      llvm::report_fatal_error("SYCL kernel name conflict");
+    assert(declaresSameEntity(FD, IT->second.getKernelEntryPointDecl()) &&
+           "SYCL kernel name conflict");
   } else {
     SYCLKernels.insert(std::make_pair(KernelNameType,
                                       BuildSYCLKernelInfo(KernelNameType, FD)));
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index cc5ef7b1502e79..2ce868c403c693 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -12053,8 +12053,7 @@ bool Sema::CheckFunctionDeclaration(Scope *S, FunctionDecl *NewFD,
   if (LangOpts.OpenMP)
     OpenMP().ActOnFinishedFunctionDefinitionInOpenMPAssumeScope(NewFD);
 
-  if (LangOpts.isSYCL() && NewFD->hasAttr<SYCLKernelEntryPointAttr>() &&
-      !NewFD->isInvalidDecl() && !NewFD->isTemplated())
+  if (LangOpts.isSYCL() && NewFD->hasAttr<SYCLKernelEntryPointAttr>())
     getASTContext().registerSYCLEntryPointFunction(NewFD);
 
   // Semantic checking for this function declaration (in isolation).
diff --git a/clang/lib/Serialization/ASTReaderDecl.cpp b/clang/lib/Serialization/ASTReaderDecl.cpp
index 7e5c08b41d5847..c8a59e9b929917 100644
--- a/clang/lib/Serialization/ASTReaderDecl.cpp
+++ b/clang/lib/Serialization/ASTReaderDecl.cpp
@@ -1159,8 +1159,7 @@ void ASTDeclReader::VisitFunctionDecl(FunctionDecl *FD) {
   // 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->isTemplated() &&
-      FD->hasAttr<SYCLKernelEntryPointAttr>()) {
+  if (FD->hasAttr<SYCLKernelEntryPointAttr>()) {
     ASTContext &C = Reader.getContext();
     C.registerSYCLEntryPointFunction(FD);
   }

From 4ecf133f5ac4858166ed04d75744bed1590fc4f4 Mon Sep 17 00:00:00 2001
From: Tom Honermann <tom.honermann@intel.com>
Date: Mon, 4 Nov 2024 07:12:56 -0800
Subject: [PATCH 9/9] Round five of addressing code review feedback from Erich
 Keane.

Changes include:
- Minor refactoring to eliminate a branch that only contained an assertion.
---
 clang/lib/AST/ASTContext.cpp | 12 +++++-------
 1 file changed, 5 insertions(+), 7 deletions(-)

diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index b7f933e6e2f746..846a1ab45a9fd0 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -14315,13 +14315,11 @@ void ASTContext::registerSYCLEntryPointFunction(FunctionDecl *FD) {
   // conflicting kernel names prior to calling this function.
   CanQualType KernelNameType = getCanonicalType(SKEPAttr->getKernelName());
   auto IT = SYCLKernels.find(KernelNameType);
-  if (IT != SYCLKernels.end()) {
-    assert(declaresSameEntity(FD, IT->second.getKernelEntryPointDecl()) &&
-           "SYCL kernel name conflict");
-  } else {
-    SYCLKernels.insert(std::make_pair(KernelNameType,
-                                      BuildSYCLKernelInfo(KernelNameType, FD)));
-  }
+  assert((IT == SYCLKernels.end() ||
+          declaresSameEntity(FD, IT->second.getKernelEntryPointDecl())) &&
+         "SYCL kernel name conflict");
+  SYCLKernels.insert(
+      std::make_pair(KernelNameType, BuildSYCLKernelInfo(KernelNameType, FD)));
 }
 
 OMPTraitInfo &ASTContext::getNewOMPTraitInfo() {