From 97159889bfe49e865dcb1248c8f6a0c4709f494b Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews@intel.com>
Date: Fri, 8 May 2020 10:53:02 -0700
Subject: [PATCH 01/14] Add support for kernel name types templated using
 enums.

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>
---
 clang/lib/Sema/SemaSYCL.cpp                   | 138 +++++++++++++++++-
 clang/test/CodeGenSYCL/int_header1.cpp        |  16 +-
 clang/test/CodeGenSYCL/integration_header.cpp |   6 +-
 .../CodeGenSYCL/kernel_name_with_typedefs.cpp |  30 ++--
 4 files changed, 161 insertions(+), 29 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 30537b847b015..bdfdc33c62fa6 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -1637,6 +1637,32 @@ static std::string eraseAnonNamespace(std::string S) {
   return S;
 }
 
+static bool checkIfDeclInCLNameSpace(DeclContext *DC) {
+
+  bool inCL = false;
+
+  while (DC) {
+
+    auto *NS = dyn_cast_or_null<NamespaceDecl>(DC);
+    DC = NS->getDeclContext();
+
+    if (NS && DC->isTranslationUnit()) {
+
+      const IdentifierInfo *II = NS->getIdentifier();
+
+      if (II && II->isStr("cl")) {
+        inCL = true;
+        break;
+      }
+    }
+
+    if (!NS)
+      break;
+  }
+
+  return inCL;
+}
+
 // Emits a forward declaration
 void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
                                         SourceLocation KernelLocation) {
@@ -1693,7 +1719,13 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
   std::string S;
   llvm::raw_string_ostream SO(S);
   D->print(SO, P);
-  O << SO.str() << ";\n";
+  O << SO.str();
+
+  if (const auto *ED = dyn_cast<EnumDecl>(D)) {
+    QualType T = ED->getIntegerType();
+    O << " : " << T.getAsString() << ";\n";
+  } else
+    O << ";\n";
 
   // print closing braces for namespaces if needed
   for (unsigned I = 0; I < NamespaceCnt; ++I)
@@ -1796,6 +1828,26 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
         }
         break;
       }
+      case TemplateArgument::ArgKind::Integral: {
+        // Handle Kernel Name Type templated using enum.
+        QualType T = Arg.getIntegralType();
+        if (const EnumType *ET = T->getAs<EnumType>()) {
+          auto EnumDecl = ET->getDecl();
+          DeclContext *DC = EnumDecl->getDeclContext();
+          auto *NS = dyn_cast_or_null<NamespaceDecl>(DC);
+          bool InCLNameSpace = false;
+
+          // Header files are included before integration header. Enums
+          // therefore defined in headers need not be forward declared. Their
+          // definitions are available in integration header.
+          if (NS)
+            InCLNameSpace = checkIfDeclInCLNameSpace(DC);
+
+          if (!InCLNameSpace)
+            emitFwdDecl(O, EnumDecl, KernelLocation);
+        }
+        break;
+      }
       default:
         break; // nop
       }
@@ -1822,6 +1874,85 @@ static std::string getCPPTypeString(QualType Ty) {
   return eraseAnonNamespace(Ty.getAsString(P));
 }
 
+static void printArguments(raw_ostream &ArgOS, ArrayRef<TemplateArgument> Args,
+                           const PrintingPolicy &P, bool ParameterPack) {
+
+  if (!ParameterPack)
+    ArgOS << "<";
+
+  bool FirstArg = true;
+  const char *Comma = ", ";
+
+  for (unsigned I = 0; I < Args.size(); I++) {
+    const TemplateArgument &Arg = Args[I];
+    if (Arg.getKind() == TemplateArgument::ArgKind::Pack) {
+      if (Arg.pack_size() && !FirstArg)
+        ArgOS << Comma;
+      printArguments(ArgOS, Arg.getPackAsArray(), P, true);
+    } else if (Arg.getKind() == TemplateArgument::ArgKind::Integral) {
+      if (!FirstArg)
+        ArgOS << Comma;
+
+      QualType T = Arg.getIntegralType();
+      const EnumType *ET = T->getAs<EnumType>();
+
+      // Enums defined in SYCL headers are not changed to explicit values since
+      // the definition is visible in integration header.
+      if (ET && !checkIfDeclInCLNameSpace(ET->getDecl()->getDeclContext())) {
+        const llvm::APSInt &Val = Arg.getAsIntegral();
+        ArgOS << "<(" << ET->getDecl()->getQualifiedNameAsString() << ")>"
+              << Val;
+      } else {
+        Arg.print(P, ArgOS);
+      }
+    } else {
+      if (!FirstArg)
+        ArgOS << Comma;
+      LangOptions LO;
+      PrintingPolicy TypePolicy(LO);
+      TypePolicy.SuppressTypedefs = true;
+      TypePolicy.SuppressTagKeyword = true;
+      Arg.print(TypePolicy, ArgOS);
+    }
+    FirstArg = false;
+  }
+
+  if (!ParameterPack)
+    ArgOS << ">";
+}
+
+static std::string getKernelNameTypeString(QualType T) {
+
+  const CXXRecordDecl *RD = T->getAsCXXRecordDecl();
+
+  if (!RD)
+    return getCPPTypeString(T);
+
+  // If kernel name type is a template specialization with enum type
+  // template parameters, enumerators in name type string should be
+  // replaced  with their underlying value since the enum definition
+  // is not visible in integration header.
+  if (const auto *TSD = dyn_cast<ClassTemplateSpecializationDecl>(RD)) {
+    LangOptions LO;
+    PrintingPolicy P(LO);
+    P.SuppressTypedefs = true;
+    SmallString<64> Buf;
+    llvm::raw_svector_ostream ArgOS(Buf);
+
+    // Print template class name
+    TSD->printQualifiedName(ArgOS, P);
+
+    const TemplateArgumentList &Args = TSD->getTemplateArgs();
+
+    // Print template arguments substituting enumerators
+    printArguments(ArgOS, Args.asArray(), P, false);
+
+    return eraseAnonNamespace(ArgOS.str().str());
+  }
+
+  return getCPPTypeString(T);
+}
+
 void SYCLIntegrationHeader::emit(raw_ostream &O) {
   O << "// This is auto-generated SYCL integration header.\n";
   O << "\n";
@@ -1938,8 +2069,9 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
         O << "', '" << c;
       O << "'> {\n";
     } else {
-      O << "template <> struct KernelInfo<" << getCPPTypeString(K.NameType)
-        << "> {\n";
+
+      O << "template <> struct KernelInfo<"
+        << getKernelNameTypeString(K.NameType) << "> {\n";
     }
     O << "  DLL_LOCAL\n";
     O << "  static constexpr const char* getName() { return \"" << K.Name
diff --git a/clang/test/CodeGenSYCL/int_header1.cpp b/clang/test/CodeGenSYCL/int_header1.cpp
index 646373e7f6adc..949fbb7cb93e6 100644
--- a/clang/test/CodeGenSYCL/int_header1.cpp
+++ b/clang/test/CodeGenSYCL/int_header1.cpp
@@ -4,14 +4,14 @@
 // CHECK:template <> struct KernelInfo<class KernelName> {
 // CHECK:template <> struct KernelInfo<::nm1::nm2::KernelName0> {
 // CHECK:template <> struct KernelInfo<::nm1::KernelName1> {
-// CHECK:template <> struct KernelInfo<::nm1::KernelName3< ::nm1::nm2::KernelName0>> {
-// CHECK:template <> struct KernelInfo<::nm1::KernelName3< ::nm1::KernelName1>> {
-// CHECK:template <> struct KernelInfo<::nm1::KernelName4< ::nm1::nm2::KernelName0>> {
-// CHECK:template <> struct KernelInfo<::nm1::KernelName4< ::nm1::KernelName1>> {
-// CHECK:template <> struct KernelInfo<::nm1::KernelName3<KernelName5>> {
-// CHECK:template <> struct KernelInfo<::nm1::KernelName4<KernelName7>> {
-// CHECK:template <> struct KernelInfo<::nm1::KernelName8< ::nm1::nm2::C>> {
-// CHECK:template <> struct KernelInfo<class TmplClassInAnonNS<class ClassInAnonNS>> {
+// CHECK:template <> struct KernelInfo<nm1::KernelName3<nm1::nm2::KernelName0>> {
+// CHECK:template <> struct KernelInfo<nm1::KernelName3<nm1::KernelName1>> {
+// CHECK:template <> struct KernelInfo<nm1::KernelName4<nm1::nm2::KernelName0>> {
+// CHECK:template <> struct KernelInfo<nm1::KernelName4<nm1::KernelName1>> {
+// CHECK:template <> struct KernelInfo<nm1::KernelName3<KernelName5>> {
+// CHECK:template <> struct KernelInfo<nm1::KernelName4<KernelName7>> {
+// CHECK:template <> struct KernelInfo<nm1::KernelName8<::nm1::nm2::C>> {
+// CHECK:template <> struct KernelInfo<TmplClassInAnonNS<ClassInAnonNS>> {
 
 // This test checks if the SYCL device compiler is able to generate correct
 // integration header when the kernel name class is expressed in different
diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp
index 84b35578f48e6..f0c2d5a04d50a 100644
--- a/clang/test/CodeGenSYCL/integration_header.cpp
+++ b/clang/test/CodeGenSYCL/integration_header.cpp
@@ -56,9 +56,9 @@
 // CHECK-NEXT: };
 //
 // CHECK: template <> struct KernelInfo<class first_kernel> {
-// CHECK: template <> struct KernelInfo<::second_namespace::second_kernel<char>> {
-// CHECK: template <> struct KernelInfo<::third_kernel<1, int, ::point<X> >> {
-// CHECK: template <> struct KernelInfo<::fourth_kernel< ::template_arg_ns::namespaced_arg<1> >> {
+// CHECK: template <> struct KernelInfo<second_namespace::second_kernel<char>> {
+// CHECK: template <> struct KernelInfo<third_kernel<1, int, point<X>>> {
+// CHECK: template <> struct KernelInfo<fourth_kernel<::template_arg_ns::namespaced_arg<1>>> {
 
 #include "sycl.hpp"
 
diff --git a/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp b/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp
index c124f95770b8f..ac54a4ecfdf1d 100644
--- a/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp
+++ b/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp
@@ -100,37 +100,37 @@ struct kernel_name4;
 int main() {
   dummy_functor f;
   // non-type template arguments
-  // CHECK: template <> struct KernelInfo<::kernel_name1<1, 1>> {
+  // CHECK: template <> struct KernelInfo<kernel_name1<1, 1>> {
   single_task<kernel_name1<1, 1>>(f);
-  // CHECK: template <> struct KernelInfo<::kernel_name1v1<1, 1>> {
+  // CHECK: template <> struct KernelInfo<kernel_name1v1<1, 1>> {
   single_task<kernel_name1v1<1, 1>>(f);
-  // CHECK: template <> struct KernelInfo<::kernel_name1v2<1, 1>> {
+  // CHECK: template <> struct KernelInfo<kernel_name1v2<1, 1>> {
   single_task<kernel_name1v2<1, 1>>(f);
   // partial template specialization
-  // CHECK: template <> struct KernelInfo<::kernel_name2<int, int>> {
+  // CHECK: template <> struct KernelInfo<kernel_name2<int, int>> {
   single_task<kernel_name2<int_t, int>>(f);
-  // CHECK: template <> struct KernelInfo<::kernel_name2<const int, char>> {
+  // CHECK: template <> struct KernelInfo<kernel_name2<const int, char>> {
   single_task<kernel_name2<cint_t, char>>(f);
-  // CHECK: template <> struct KernelInfo<::kernel_name2<long, float>> {
+  // CHECK: template <> struct KernelInfo<kernel_name2<long, float>> {
   single_task<kernel_name2<space::long_t, float>>(f);
-  // CHECK: template <> struct KernelInfo<::kernel_name2<const long, ::A>> {
+  // CHECK: template <> struct KernelInfo<kernel_name2<const long, A>> {
   single_task<kernel_name2<const space::long_t, space::a_t>>(f);
-  // CHECK: template <> struct KernelInfo<::kernel_name2<const volatile long, const ::space::B>> {
+  // CHECK: template <> struct KernelInfo<kernel_name2<const volatile long, const space::B>> {
   single_task<kernel_name2<volatile space::clong_t, const space::b_t>>(f);
-  // CHECK: template <> struct KernelInfo<::kernel_name2< ::A, long>> {
+  // CHECK: template <> struct KernelInfo<kernel_name2<A, long>> {
   single_task<kernel_name2<space::a_t, space::long_t>>(f);
-  // CHECK: template <> struct KernelInfo<::kernel_name2< ::space::B, int>> {
+  // CHECK: template <> struct KernelInfo<kernel_name2<space::B, int>> {
   single_task<kernel_name2<space::b_t, int_t>>(f);
   // full template specialization
-  // CHECK: template <> struct KernelInfo<::kernel_name2<int, const unsigned int>> {
+  // CHECK: template <> struct KernelInfo<kernel_name2<int, const unsigned int>> {
   single_task<kernel_name2<int_t, const uint_t>>(f);
-  // CHECK: template <> struct KernelInfo<::kernel_name2<const long, volatile const unsigned long>> {
+  // CHECK: template <> struct KernelInfo<kernel_name2<const long, const volatile unsigned long>> {
   single_task<kernel_name2<space::clong_t, volatile space::culong_t>>(f);
-  // CHECK: template <> struct KernelInfo<::kernel_name2< ::A, volatile ::space::B>> {
+  // CHECK: template <> struct KernelInfo<kernel_name2<A, volatile space::B>> {
   single_task<kernel_name2<space::a_t, volatile space::b_t>>(f);
-  // CHECK: template <> struct KernelInfo<::kernel_name3<1>> {
+  // CHECK: template <> struct KernelInfo<kernel_name3<1>> {
   single_task<kernel_name3<1>>(f);
-  // CHECK: template <> struct KernelInfo<::kernel_name4<1>> {
+  // CHECK: template <> struct KernelInfo<kernel_name4<1>> {
   single_task<kernel_name4<1>>(f);
 
   return 0;

From b337e438cb2697595fb5df5afb27854093a74d8b Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews@intel.com>
Date: Sun, 10 May 2020 21:35:56 -0700
Subject: [PATCH 02/14] Fix crash of enum in TU scope. Add test

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>
---
 clang/lib/Sema/SemaSYCL.cpp                | 22 +++----
 clang/test/CodeGenSYCL/kernelname-enum.cpp | 77 ++++++++++++++++++++++
 2 files changed, 85 insertions(+), 14 deletions(-)
 create mode 100644 clang/test/CodeGenSYCL/kernelname-enum.cpp

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index bdfdc33c62fa6..5553e3a3e6100 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -1642,24 +1642,19 @@ static bool checkIfDeclInCLNameSpace(DeclContext *DC) {
   bool inCL = false;
 
   while (DC) {
-
     auto *NS = dyn_cast_or_null<NamespaceDecl>(DC);
-    DC = NS->getDeclContext();
 
-    if (NS && DC->isTranslationUnit()) {
+    if (!NS)
+      break;
 
+    DC = NS->getDeclContext();
+    if (DC->isTranslationUnit()) {
       const IdentifierInfo *II = NS->getIdentifier();
-
-      if (II && II->isStr("cl")) {
+      if (II && II->isStr("cl"))
         inCL = true;
-        break;
-      }
-    }
-
-    if (!NS)
       break;
+    }
   }
-
   return inCL;
 }
 
@@ -1843,7 +1838,7 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
           if (NS)
             InCLNameSpace = checkIfDeclInCLNameSpace(DC);
 
-          if (!InCLNameSpace)
+          if (!NS || (NS && !InCLNameSpace))
             emitFwdDecl(O, EnumDecl, KernelLocation);
         }
         break;
@@ -1900,8 +1895,7 @@ static void printArguments(raw_ostream &ArgOS, ArrayRef<TemplateArgument> Args,
       // the definition is visible in integration header.
       if (ET && !checkIfDeclInCLNameSpace(ET->getDecl()->getDeclContext())) {
         const llvm::APSInt &Val = Arg.getAsIntegral();
-        ArgOS << "<(" << ET->getDecl()->getQualifiedNameAsString() << ")>"
-              << Val;
+        ArgOS << "(" << ET->getDecl()->getQualifiedNameAsString() << ")" << Val;
       } else {
         Arg.print(P, ArgOS);
       }
diff --git a/clang/test/CodeGenSYCL/kernelname-enum.cpp b/clang/test/CodeGenSYCL/kernelname-enum.cpp
new file mode 100644
index 0000000000000..8b02f97635e37
--- /dev/null
+++ b/clang/test/CodeGenSYCL/kernelname-enum.cpp
@@ -0,0 +1,77 @@
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -fsyntax-only
+// RUN: FileCheck -input-file=%t.h %s
+
+#include "sycl.hpp"
+
+enum class no_namespace_int : int {
+  val_1,
+  val_2
+};
+
+enum class no_namespace_short : short {
+  val_1,
+  val_2
+};
+
+namespace internal {
+enum class namespace_short : short {
+  val_1,
+  val_2
+};
+}
+
+template <no_namespace_int EnumType>
+class dummy_functor_1 {
+public:
+  void operator()() {}
+};
+
+template <no_namespace_short EnumType>
+class dummy_functor_2 {
+public:
+  void operator()() {}
+};
+
+template <internal::namespace_short EnumType>
+class dummy_functor_3 {
+public:
+  void operator()() {}
+};
+
+int main() {
+
+  dummy_functor_1<no_namespace_int::val_1> f1;
+  dummy_functor_2<no_namespace_short::val_2> f2;
+  dummy_functor_3<internal::namespace_short::val_2> f3;
+
+  cl::sycl::queue q;
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f1);
+  });
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f2);
+  });
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f3);
+  });
+
+  return 0;
+}
+
+// CHECK: Forward declarations of templated kernel function types:
+// CHECK: enum class no_namespace_int : int;
+// CHECK: template <no_namespace_int EnumType> class dummy_functor_1;
+// CHECK: enum class no_namespace_short : short;
+// CHECK: template <no_namespace_short EnumType> class dummy_functor_2;
+// CHECK: namespace internal {
+// CHECK-NEXT: enum class namespace_short : short;
+// CHECK-NEXT: }
+// CHECK: template <internal::namespace_short EnumType> class dummy_functor_3;
+
+// CHECK: Specializations of KernelInfo for kernel function types:
+// CHECK: template <> struct KernelInfo<dummy_functor_1<(no_namespace_int)0>>
+// CHECK: template <> struct KernelInfo<dummy_functor_2<(no_namespace_short)1>>
+// CHECK: template <> struct KernelInfo<dummy_functor_3<(internal::namespace_short)1>>

From c182c488d5df4ab38f030bf55eeb692271234a4c Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews@intel.com>
Date: Mon, 11 May 2020 22:59:24 -0700
Subject: [PATCH 03/14] Modify printing policy and add test for enums in
 anonymous namespace

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>
---
 clang/lib/Sema/SemaSYCL.cpp                |  1 +
 clang/test/CodeGenSYCL/kernelname-enum.cpp | 23 ++++++++++++++++++++++
 2 files changed, 24 insertions(+)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 5553e3a3e6100..70bb5aa6321ad 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -1711,6 +1711,7 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
   PrintingPolicy P(D->getASTContext().getLangOpts());
   P.adjustForCPlusPlusFwdDecl();
   P.SuppressTypedefs = true;
+  P.SuppressUnwrittenScope = true;
   std::string S;
   llvm::raw_string_ostream SO(S);
   D->print(SO, P);
diff --git a/clang/test/CodeGenSYCL/kernelname-enum.cpp b/clang/test/CodeGenSYCL/kernelname-enum.cpp
index 8b02f97635e37..8da0d224626f6 100644
--- a/clang/test/CodeGenSYCL/kernelname-enum.cpp
+++ b/clang/test/CodeGenSYCL/kernelname-enum.cpp
@@ -20,6 +20,13 @@ enum class namespace_short : short {
 };
 }
 
+namespace {
+enum class enum_in_anonNS : short {
+  val_1,
+  val_2
+};
+}
+
 template <no_namespace_int EnumType>
 class dummy_functor_1 {
 public:
@@ -38,11 +45,18 @@ class dummy_functor_3 {
   void operator()() {}
 };
 
+template <enum_in_anonNS EnumType>
+class dummy_functor_4 {
+public:
+  void operator()() {}
+};
+
 int main() {
 
   dummy_functor_1<no_namespace_int::val_1> f1;
   dummy_functor_2<no_namespace_short::val_2> f2;
   dummy_functor_3<internal::namespace_short::val_2> f3;
+  dummy_functor_4<enum_in_anonNS::val_2> f4;
 
   cl::sycl::queue q;
 
@@ -58,6 +72,10 @@ int main() {
     cgh.single_task(f3);
   });
 
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f4);
+  });
+
   return 0;
 }
 
@@ -70,8 +88,13 @@ int main() {
 // CHECK-NEXT: enum class namespace_short : short;
 // CHECK-NEXT: }
 // CHECK: template <internal::namespace_short EnumType> class dummy_functor_3;
+// CHECK: namespace  {
+// CHECK-NEXT: enum class enum_in_anonNS : short;
+// CHECK-NEXT: }
+// CHECK: template <enum_in_anonNS EnumType> class dummy_functor_4;
 
 // CHECK: Specializations of KernelInfo for kernel function types:
 // CHECK: template <> struct KernelInfo<dummy_functor_1<(no_namespace_int)0>>
 // CHECK: template <> struct KernelInfo<dummy_functor_2<(no_namespace_short)1>>
 // CHECK: template <> struct KernelInfo<dummy_functor_3<(internal::namespace_short)1>>
+// CHECK: template <> struct KernelInfo<dummy_functor_4<(enum_in_anonNS)1>>

From f6b9b590bc20002c2bef5b915bc61b238c91f18d Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews@intel.com>
Date: Thu, 14 May 2020 06:29:48 -0700
Subject: [PATCH 04/14] Implemented review comments:  Removed check for CL
 namespace.  Emit scope resolution operator for global namespace.  Added test
 when underlying type is not fixed.

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>
---
 clang/include/clang/AST/Decl.h                |  7 +-
 clang/lib/AST/Decl.cpp                        | 18 ++++--
 clang/lib/Sema/SemaSYCL.cpp                   | 64 ++++++-------------
 clang/test/CodeGenSYCL/int_header1.cpp        | 14 ++--
 clang/test/CodeGenSYCL/integration_header.cpp |  6 +-
 .../CodeGenSYCL/kernel_name_with_typedefs.cpp | 30 ++++-----
 clang/test/CodeGenSYCL/kernelname-enum.cpp    | 27 ++++++--
 7 files changed, 84 insertions(+), 82 deletions(-)

diff --git a/clang/include/clang/AST/Decl.h b/clang/include/clang/AST/Decl.h
index 7db74e0803cec..9940636192f1d 100644
--- a/clang/include/clang/AST/Decl.h
+++ b/clang/include/clang/AST/Decl.h
@@ -283,15 +283,16 @@ class NamedDecl : public Decl {
   /// Creating this name is expensive, so it should be called only when
   /// performance doesn't matter.
   void printQualifiedName(raw_ostream &OS) const;
-  void printQualifiedName(raw_ostream &OS, const PrintingPolicy &Policy) const;
+  void printQualifiedName(raw_ostream &OS, const PrintingPolicy &Policy,
+                          bool WithGlobalNsPrefix = false) const;
 
   /// Print only the nested name specifier part of a fully-qualified name,
   /// including the '::' at the end. E.g.
   ///    when `printQualifiedName(D)` prints "A::B::i",
   ///    this function prints "A::B::".
   void printNestedNameSpecifier(raw_ostream &OS) const;
-  void printNestedNameSpecifier(raw_ostream &OS,
-                                const PrintingPolicy &Policy) const;
+  void printNestedNameSpecifier(raw_ostream &OS, const PrintingPolicy &Policy,
+                                bool WithGlobalNsPrefix = false) const;
 
   // FIXME: Remove string version.
   std::string getQualifiedNameAsString() const;
diff --git a/clang/lib/AST/Decl.cpp b/clang/lib/AST/Decl.cpp
index 9c1b99d30e788..24a7dc38d1880 100644
--- a/clang/lib/AST/Decl.cpp
+++ b/clang/lib/AST/Decl.cpp
@@ -1538,14 +1538,14 @@ void NamedDecl::printQualifiedName(raw_ostream &OS) const {
   printQualifiedName(OS, getASTContext().getPrintingPolicy());
 }
 
-void NamedDecl::printQualifiedName(raw_ostream &OS,
-                                   const PrintingPolicy &P) const {
+void NamedDecl::printQualifiedName(raw_ostream &OS, const PrintingPolicy &P,
+                                   bool WithGlobalNsPrefix) const {
   if (getDeclContext()->isFunctionOrMethod()) {
     // We do not print '(anonymous)' for function parameters without name.
     printName(OS);
     return;
   }
-  printNestedNameSpecifier(OS, P);
+  printNestedNameSpecifier(OS, P, WithGlobalNsPrefix);
   if (getDeclName())
     OS << *this;
   else {
@@ -1566,7 +1566,8 @@ void NamedDecl::printNestedNameSpecifier(raw_ostream &OS) const {
 }
 
 void NamedDecl::printNestedNameSpecifier(raw_ostream &OS,
-                                         const PrintingPolicy &P) const {
+                                         const PrintingPolicy &P,
+                                         bool WithGlobalNsPrefix) const {
   const DeclContext *Ctx = getDeclContext();
 
   // For ObjC methods and properties, look through categories and use the
@@ -1593,6 +1594,11 @@ void NamedDecl::printNestedNameSpecifier(raw_ostream &OS,
     Ctx = Ctx->getParent();
   }
 
+  if (Contexts.empty() && WithGlobalNsPrefix) {
+    if (getDeclContext()->isTranslationUnit())
+      OS << "::";
+  }
+
   for (const DeclContext *DC : llvm::reverse(Contexts)) {
     if (const auto *Spec = dyn_cast<ClassTemplateSpecializationDecl>(DC)) {
       OS << Spec->getName();
@@ -1605,7 +1611,9 @@ void NamedDecl::printNestedNameSpecifier(raw_ostream &OS,
       if (ND->isAnonymousNamespace()) {
         OS << (P.MSVCFormatting ? "`anonymous namespace\'"
                                 : "(anonymous namespace)");
-      }
+      } else if (WithGlobalNsPrefix &&
+                 ND->getDeclContext()->isTranslationUnit())
+        OS << "::" << *ND;
       else
         OS << *ND;
     } else if (const auto *RD = dyn_cast<RecordDecl>(DC)) {
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 70bb5aa6321ad..d7ac1458a5c5f 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -1637,27 +1637,6 @@ static std::string eraseAnonNamespace(std::string S) {
   return S;
 }
 
-static bool checkIfDeclInCLNameSpace(DeclContext *DC) {
-
-  bool inCL = false;
-
-  while (DC) {
-    auto *NS = dyn_cast_or_null<NamespaceDecl>(DC);
-
-    if (!NS)
-      break;
-
-    DC = NS->getDeclContext();
-    if (DC->isTranslationUnit()) {
-      const IdentifierInfo *II = NS->getIdentifier();
-      if (II && II->isStr("cl"))
-        inCL = true;
-      break;
-    }
-  }
-  return inCL;
-}
-
 // Emits a forward declaration
 void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
                                         SourceLocation KernelLocation) {
@@ -1828,19 +1807,8 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
         // Handle Kernel Name Type templated using enum.
         QualType T = Arg.getIntegralType();
         if (const EnumType *ET = T->getAs<EnumType>()) {
-          auto EnumDecl = ET->getDecl();
-          DeclContext *DC = EnumDecl->getDeclContext();
-          auto *NS = dyn_cast_or_null<NamespaceDecl>(DC);
-          bool InCLNameSpace = false;
-
-          // Header files are included before integration header. Enums
-          // therefore defined in headers need not be forward declared. Their
-          // definitions are available in integration header.
-          if (NS)
-            InCLNameSpace = checkIfDeclInCLNameSpace(DC);
-
-          if (!NS || (NS && !InCLNameSpace))
-            emitFwdDecl(O, EnumDecl, KernelLocation);
+          const EnumDecl *ED = ET->getDecl();
+          emitFwdDecl(O, ED, KernelLocation);
         }
         break;
       }
@@ -1870,7 +1838,8 @@ static std::string getCPPTypeString(QualType Ty) {
   return eraseAnonNamespace(Ty.getAsString(P));
 }
 
-static void printArguments(raw_ostream &ArgOS, ArrayRef<TemplateArgument> Args,
+static void printArguments(ASTContext &Ctx, raw_ostream &ArgOS,
+                           ArrayRef<TemplateArgument> Args,
                            const PrintingPolicy &P, bool ParameterPack) {
 
   if (!ParameterPack)
@@ -1884,7 +1853,7 @@ static void printArguments(raw_ostream &ArgOS, ArrayRef<TemplateArgument> Args,
     if (Arg.getKind() == TemplateArgument::ArgKind::Pack) {
       if (Arg.pack_size() && !FirstArg)
         ArgOS << Comma;
-      printArguments(ArgOS, Arg.getPackAsArray(), P, true);
+      printArguments(Ctx, ArgOS, Arg.getPackAsArray(), P, true);
     } else if (Arg.getKind() == TemplateArgument::ArgKind::Integral) {
       if (!FirstArg)
         ArgOS << Comma;
@@ -1892,22 +1861,27 @@ static void printArguments(raw_ostream &ArgOS, ArrayRef<TemplateArgument> Args,
       QualType T = Arg.getIntegralType();
       const EnumType *ET = T->getAs<EnumType>();
 
-      // Enums defined in SYCL headers are not changed to explicit values since
-      // the definition is visible in integration header.
-      if (ET && !checkIfDeclInCLNameSpace(ET->getDecl()->getDeclContext())) {
+      if (ET) {
         const llvm::APSInt &Val = Arg.getAsIntegral();
         ArgOS << "(" << ET->getDecl()->getQualifiedNameAsString() << ")" << Val;
       } else {
         Arg.print(P, ArgOS);
       }
-    } else {
+    } else if (Arg.getKind() == TemplateArgument::ArgKind::Type) {
       if (!FirstArg)
         ArgOS << Comma;
       LangOptions LO;
       PrintingPolicy TypePolicy(LO);
       TypePolicy.SuppressTypedefs = true;
       TypePolicy.SuppressTagKeyword = true;
-      Arg.print(TypePolicy, ArgOS);
+      QualType T = Arg.getAsType();
+      QualType FullyQualifiedType =
+          TypeName::getFullyQualifiedType(T, Ctx, true);
+      ArgOS << FullyQualifiedType.getAsString(TypePolicy);
+    } else {
+      if (!FirstArg)
+        ArgOS << Comma;
+      Arg.print(P, ArgOS);
     }
     FirstArg = false;
   }
@@ -1935,12 +1909,12 @@ static std::string getKernelNameTypeString(QualType T) {
     llvm::raw_svector_ostream ArgOS(Buf);
 
     // Print template class name
-    TSD->printQualifiedName(ArgOS, P);
-
-    const TemplateArgumentList &Args = TSD->getTemplateArgs();
+    TSD->printQualifiedName(ArgOS, P, true);
 
     // Print template arguments substituting enumerators
-    printArguments(ArgOS, Args.asArray(), P, false);
+    ASTContext &Ctx = RD->getASTContext();
+    const TemplateArgumentList &Args = TSD->getTemplateArgs();
+    printArguments(Ctx, ArgOS, Args.asArray(), P, false);
 
     return eraseAnonNamespace(ArgOS.str().str());
   }
diff --git a/clang/test/CodeGenSYCL/int_header1.cpp b/clang/test/CodeGenSYCL/int_header1.cpp
index 949fbb7cb93e6..1b7baac056e5e 100644
--- a/clang/test/CodeGenSYCL/int_header1.cpp
+++ b/clang/test/CodeGenSYCL/int_header1.cpp
@@ -4,13 +4,13 @@
 // CHECK:template <> struct KernelInfo<class KernelName> {
 // CHECK:template <> struct KernelInfo<::nm1::nm2::KernelName0> {
 // CHECK:template <> struct KernelInfo<::nm1::KernelName1> {
-// CHECK:template <> struct KernelInfo<nm1::KernelName3<nm1::nm2::KernelName0>> {
-// CHECK:template <> struct KernelInfo<nm1::KernelName3<nm1::KernelName1>> {
-// CHECK:template <> struct KernelInfo<nm1::KernelName4<nm1::nm2::KernelName0>> {
-// CHECK:template <> struct KernelInfo<nm1::KernelName4<nm1::KernelName1>> {
-// CHECK:template <> struct KernelInfo<nm1::KernelName3<KernelName5>> {
-// CHECK:template <> struct KernelInfo<nm1::KernelName4<KernelName7>> {
-// CHECK:template <> struct KernelInfo<nm1::KernelName8<::nm1::nm2::C>> {
+// CHECK:template <> struct KernelInfo<::nm1::KernelName3<::nm1::nm2::KernelName0>> {
+// CHECK:template <> struct KernelInfo<::nm1::KernelName3<::nm1::KernelName1>> {
+// CHECK:template <> struct KernelInfo<::nm1::KernelName4<::nm1::nm2::KernelName0>> {
+// CHECK:template <> struct KernelInfo<::nm1::KernelName4<::nm1::KernelName1>> {
+// CHECK:template <> struct KernelInfo<::nm1::KernelName3<KernelName5>> {
+// CHECK:template <> struct KernelInfo<::nm1::KernelName4<KernelName7>> {
+// CHECK:template <> struct KernelInfo<::nm1::KernelName8<::nm1::nm2::C>> {
 // CHECK:template <> struct KernelInfo<TmplClassInAnonNS<ClassInAnonNS>> {
 
 // This test checks if the SYCL device compiler is able to generate correct
diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp
index f0c2d5a04d50a..58d0c3addcd8c 100644
--- a/clang/test/CodeGenSYCL/integration_header.cpp
+++ b/clang/test/CodeGenSYCL/integration_header.cpp
@@ -56,9 +56,9 @@
 // CHECK-NEXT: };
 //
 // CHECK: template <> struct KernelInfo<class first_kernel> {
-// CHECK: template <> struct KernelInfo<second_namespace::second_kernel<char>> {
-// CHECK: template <> struct KernelInfo<third_kernel<1, int, point<X>>> {
-// CHECK: template <> struct KernelInfo<fourth_kernel<::template_arg_ns::namespaced_arg<1>>> {
+// CHECK: template <> struct KernelInfo<::second_namespace::second_kernel<char>> {
+// CHECK: template <> struct KernelInfo<::third_kernel<1, int, ::point<X>>> {
+// CHECK: template <> struct KernelInfo<::fourth_kernel<::template_arg_ns::namespaced_arg<1>>> {
 
 #include "sycl.hpp"
 
diff --git a/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp b/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp
index ac54a4ecfdf1d..2ad4c70020281 100644
--- a/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp
+++ b/clang/test/CodeGenSYCL/kernel_name_with_typedefs.cpp
@@ -100,37 +100,37 @@ struct kernel_name4;
 int main() {
   dummy_functor f;
   // non-type template arguments
-  // CHECK: template <> struct KernelInfo<kernel_name1<1, 1>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name1<1, 1>> {
   single_task<kernel_name1<1, 1>>(f);
-  // CHECK: template <> struct KernelInfo<kernel_name1v1<1, 1>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name1v1<1, 1>> {
   single_task<kernel_name1v1<1, 1>>(f);
-  // CHECK: template <> struct KernelInfo<kernel_name1v2<1, 1>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name1v2<1, 1>> {
   single_task<kernel_name1v2<1, 1>>(f);
   // partial template specialization
-  // CHECK: template <> struct KernelInfo<kernel_name2<int, int>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name2<int, int>> {
   single_task<kernel_name2<int_t, int>>(f);
-  // CHECK: template <> struct KernelInfo<kernel_name2<const int, char>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name2<const int, char>> {
   single_task<kernel_name2<cint_t, char>>(f);
-  // CHECK: template <> struct KernelInfo<kernel_name2<long, float>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name2<long, float>> {
   single_task<kernel_name2<space::long_t, float>>(f);
-  // CHECK: template <> struct KernelInfo<kernel_name2<const long, A>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name2<const long, ::A>> {
   single_task<kernel_name2<const space::long_t, space::a_t>>(f);
-  // CHECK: template <> struct KernelInfo<kernel_name2<const volatile long, const space::B>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name2<const volatile long, const ::space::B>> {
   single_task<kernel_name2<volatile space::clong_t, const space::b_t>>(f);
-  // CHECK: template <> struct KernelInfo<kernel_name2<A, long>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name2<::A, long>> {
   single_task<kernel_name2<space::a_t, space::long_t>>(f);
-  // CHECK: template <> struct KernelInfo<kernel_name2<space::B, int>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name2<::space::B, int>> {
   single_task<kernel_name2<space::b_t, int_t>>(f);
   // full template specialization
-  // CHECK: template <> struct KernelInfo<kernel_name2<int, const unsigned int>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name2<int, const unsigned int>> {
   single_task<kernel_name2<int_t, const uint_t>>(f);
-  // CHECK: template <> struct KernelInfo<kernel_name2<const long, const volatile unsigned long>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name2<const long, const volatile unsigned long>> {
   single_task<kernel_name2<space::clong_t, volatile space::culong_t>>(f);
-  // CHECK: template <> struct KernelInfo<kernel_name2<A, volatile space::B>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name2<::A, volatile ::space::B>> {
   single_task<kernel_name2<space::a_t, volatile space::b_t>>(f);
-  // CHECK: template <> struct KernelInfo<kernel_name3<1>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name3<1>> {
   single_task<kernel_name3<1>>(f);
-  // CHECK: template <> struct KernelInfo<kernel_name4<1>> {
+  // CHECK: template <> struct KernelInfo<::kernel_name4<1>> {
   single_task<kernel_name4<1>>(f);
 
   return 0;
diff --git a/clang/test/CodeGenSYCL/kernelname-enum.cpp b/clang/test/CodeGenSYCL/kernelname-enum.cpp
index 8da0d224626f6..dd3429a8667c8 100644
--- a/clang/test/CodeGenSYCL/kernelname-enum.cpp
+++ b/clang/test/CodeGenSYCL/kernelname-enum.cpp
@@ -27,6 +27,11 @@ enum class enum_in_anonNS : short {
 };
 }
 
+enum class no_type_set {
+  val_1,
+  val_2
+};
+
 template <no_namespace_int EnumType>
 class dummy_functor_1 {
 public:
@@ -51,12 +56,19 @@ class dummy_functor_4 {
   void operator()() {}
 };
 
+template <no_type_set EnumType>
+class dummy_functor_5 {
+public:
+  void operator()() {}
+};
+
 int main() {
 
   dummy_functor_1<no_namespace_int::val_1> f1;
   dummy_functor_2<no_namespace_short::val_2> f2;
   dummy_functor_3<internal::namespace_short::val_2> f3;
   dummy_functor_4<enum_in_anonNS::val_2> f4;
+  dummy_functor_5<no_type_set::val_1> f5;
 
   cl::sycl::queue q;
 
@@ -76,6 +88,10 @@ int main() {
     cgh.single_task(f4);
   });
 
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f5);
+  });
+
   return 0;
 }
 
@@ -92,9 +108,12 @@ int main() {
 // CHECK-NEXT: enum class enum_in_anonNS : short;
 // CHECK-NEXT: }
 // CHECK: template <enum_in_anonNS EnumType> class dummy_functor_4;
+// CHECK: enum class no_type_set : int;
+// CHECK: template <no_type_set EnumType> class dummy_functor_5;
 
 // CHECK: Specializations of KernelInfo for kernel function types:
-// CHECK: template <> struct KernelInfo<dummy_functor_1<(no_namespace_int)0>>
-// CHECK: template <> struct KernelInfo<dummy_functor_2<(no_namespace_short)1>>
-// CHECK: template <> struct KernelInfo<dummy_functor_3<(internal::namespace_short)1>>
-// CHECK: template <> struct KernelInfo<dummy_functor_4<(enum_in_anonNS)1>>
+// CHECK: template <> struct KernelInfo<::dummy_functor_1<(no_namespace_int)0>>
+// CHECK: template <> struct KernelInfo<::dummy_functor_2<(no_namespace_short)1>>
+// CHECK: template <> struct KernelInfo<::dummy_functor_3<(internal::namespace_short)1>>
+// CHECK: template <> struct KernelInfo<::dummy_functor_4<(enum_in_anonNS)1>>
+// CHECK: template <> struct KernelInfo<::dummy_functor_5<(no_type_set)0>>

From 07135abb7613f6aee28b78e60537918da30cba99 Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews@intel.com>
Date: Thu, 14 May 2020 06:51:36 -0700
Subject: [PATCH 05/14] Add back up promotion type to handle enum forward
 declarations with no fixed underlying type

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>
---
 clang/lib/Sema/SemaSYCL.cpp | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index d7ac1458a5c5f..60f0633490d21 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -1698,6 +1698,10 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
 
   if (const auto *ED = dyn_cast<EnumDecl>(D)) {
     QualType T = ED->getIntegerType();
+    // Backup since getIntegerType() returns null for enum forward
+    // declaration with no fixed underlying type
+    if (!T)
+      T = ED->getPromotionType();
     O << " : " << T.getAsString() << ";\n";
   } else
     O << ";\n";

From df0994a471e17d95e8f82c2bec75e602bb3a20bd Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews@intel.com>
Date: Thu, 14 May 2020 08:27:18 -0700
Subject: [PATCH 06/14] Implemented review comments including:

  Minor refactoring when emitting scope resolution operator.
  Added tests got unscoped enum.
  Added comment syntax for bool argument.

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>
---
 clang/lib/AST/Decl.cpp                     | 10 ++----
 clang/lib/Sema/SemaSYCL.cpp                | 18 +++++-----
 clang/test/CodeGenSYCL/int_header1.cpp     |  2 +-
 clang/test/CodeGenSYCL/kernelname-enum.cpp | 38 ++++++++++++++++++++++
 4 files changed, 52 insertions(+), 16 deletions(-)

diff --git a/clang/lib/AST/Decl.cpp b/clang/lib/AST/Decl.cpp
index 24a7dc38d1880..e286388387a09 100644
--- a/clang/lib/AST/Decl.cpp
+++ b/clang/lib/AST/Decl.cpp
@@ -1594,9 +1594,8 @@ void NamedDecl::printNestedNameSpecifier(raw_ostream &OS,
     Ctx = Ctx->getParent();
   }
 
-  if (Contexts.empty() && WithGlobalNsPrefix) {
-    if (getDeclContext()->isTranslationUnit())
-      OS << "::";
+  if (WithGlobalNsPrefix) {
+    OS << "::";
   }
 
   for (const DeclContext *DC : llvm::reverse(Contexts)) {
@@ -1611,10 +1610,7 @@ void NamedDecl::printNestedNameSpecifier(raw_ostream &OS,
       if (ND->isAnonymousNamespace()) {
         OS << (P.MSVCFormatting ? "`anonymous namespace\'"
                                 : "(anonymous namespace)");
-      } else if (WithGlobalNsPrefix &&
-                 ND->getDeclContext()->isTranslationUnit())
-        OS << "::" << *ND;
-      else
+      } else
         OS << *ND;
     } else if (const auto *RD = dyn_cast<RecordDecl>(DC)) {
       if (!RD->getIdentifier())
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 60f0633490d21..fdb8940e31e15 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -1700,11 +1700,12 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
     QualType T = ED->getIntegerType();
     // Backup since getIntegerType() returns null for enum forward
     // declaration with no fixed underlying type
-    if (!T)
+    if (T.isNull())
       T = ED->getPromotionType();
-    O << " : " << T.getAsString() << ";\n";
-  } else
-    O << ";\n";
+    O << " : " << T.getAsString();
+  }
+
+  O << ";\n";
 
   // print closing braces for namespaces if needed
   for (unsigned I = 0; I < NamespaceCnt; ++I)
@@ -1810,7 +1811,7 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
       case TemplateArgument::ArgKind::Integral: {
         // Handle Kernel Name Type templated using enum.
         QualType T = Arg.getIntegralType();
-        if (const EnumType *ET = T->getAs<EnumType>()) {
+        if (const auto *ET = T->getAs<EnumType>()) {
           const EnumDecl *ED = ET->getDecl();
           emitFwdDecl(O, ED, KernelLocation);
         }
@@ -1857,7 +1858,8 @@ static void printArguments(ASTContext &Ctx, raw_ostream &ArgOS,
     if (Arg.getKind() == TemplateArgument::ArgKind::Pack) {
       if (Arg.pack_size() && !FirstArg)
         ArgOS << Comma;
-      printArguments(Ctx, ArgOS, Arg.getPackAsArray(), P, true);
+      printArguments(Ctx, ArgOS, Arg.getPackAsArray(), P,
+                     /*ParameterPack*/ true);
     } else if (Arg.getKind() == TemplateArgument::ArgKind::Integral) {
       if (!FirstArg)
         ArgOS << Comma;
@@ -1913,12 +1915,12 @@ static std::string getKernelNameTypeString(QualType T) {
     llvm::raw_svector_ostream ArgOS(Buf);
 
     // Print template class name
-    TSD->printQualifiedName(ArgOS, P, true);
+    TSD->printQualifiedName(ArgOS, P, /*WithGlobalNsPrefix*/ true);
 
     // Print template arguments substituting enumerators
     ASTContext &Ctx = RD->getASTContext();
     const TemplateArgumentList &Args = TSD->getTemplateArgs();
-    printArguments(Ctx, ArgOS, Args.asArray(), P, false);
+    printArguments(Ctx, ArgOS, Args.asArray(), P, /*ParameterPack*/ false);
 
     return eraseAnonNamespace(ArgOS.str().str());
   }
diff --git a/clang/test/CodeGenSYCL/int_header1.cpp b/clang/test/CodeGenSYCL/int_header1.cpp
index 1b7baac056e5e..a2e6288a7875c 100644
--- a/clang/test/CodeGenSYCL/int_header1.cpp
+++ b/clang/test/CodeGenSYCL/int_header1.cpp
@@ -11,7 +11,7 @@
 // CHECK:template <> struct KernelInfo<::nm1::KernelName3<KernelName5>> {
 // CHECK:template <> struct KernelInfo<::nm1::KernelName4<KernelName7>> {
 // CHECK:template <> struct KernelInfo<::nm1::KernelName8<::nm1::nm2::C>> {
-// CHECK:template <> struct KernelInfo<TmplClassInAnonNS<ClassInAnonNS>> {
+// CHECK:template <> struct KernelInfo<::TmplClassInAnonNS<ClassInAnonNS>> {
 
 // This test checks if the SYCL device compiler is able to generate correct
 // integration header when the kernel name class is expressed in different
diff --git a/clang/test/CodeGenSYCL/kernelname-enum.cpp b/clang/test/CodeGenSYCL/kernelname-enum.cpp
index dd3429a8667c8..8ae470a4f402b 100644
--- a/clang/test/CodeGenSYCL/kernelname-enum.cpp
+++ b/clang/test/CodeGenSYCL/kernelname-enum.cpp
@@ -3,6 +3,16 @@
 
 #include "sycl.hpp"
 
+enum unscoped_enum : int {
+  val_1,
+  val_2
+};
+
+enum unscoped_enum_no_type_set {
+  val_3,
+  val_4
+};
+
 enum class no_namespace_int : int {
   val_1,
   val_2
@@ -62,6 +72,18 @@ class dummy_functor_5 {
   void operator()() {}
 };
 
+template <unscoped_enum EnumType>
+class dummy_functor_6 {
+public:
+  void operator()() {}
+};
+
+template <unscoped_enum_no_type_set EnumType>
+class dummy_functor_7 {
+public:
+  void operator()() {}
+};
+
 int main() {
 
   dummy_functor_1<no_namespace_int::val_1> f1;
@@ -69,6 +91,8 @@ int main() {
   dummy_functor_3<internal::namespace_short::val_2> f3;
   dummy_functor_4<enum_in_anonNS::val_2> f4;
   dummy_functor_5<no_type_set::val_1> f5;
+  dummy_functor_6<unscoped_enum::val_1> f6;
+  dummy_functor_7<unscoped_enum_no_type_set::val_4> f7;
 
   cl::sycl::queue q;
 
@@ -92,6 +116,14 @@ int main() {
     cgh.single_task(f5);
   });
 
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f6);
+  });
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f7);
+  });
+
   return 0;
 }
 
@@ -110,6 +142,10 @@ int main() {
 // CHECK: template <enum_in_anonNS EnumType> class dummy_functor_4;
 // CHECK: enum class no_type_set : int;
 // CHECK: template <no_type_set EnumType> class dummy_functor_5;
+// CHECK: enum unscoped_enum : int;
+// CHECK: template <unscoped_enum EnumType> class dummy_functor_6;
+// CHECK: enum unscoped_enum_no_type_set : unsigned int;
+// CHECK: template <unscoped_enum_no_type_set EnumType> class dummy_functor_7;
 
 // CHECK: Specializations of KernelInfo for kernel function types:
 // CHECK: template <> struct KernelInfo<::dummy_functor_1<(no_namespace_int)0>>
@@ -117,3 +153,5 @@ int main() {
 // CHECK: template <> struct KernelInfo<::dummy_functor_3<(internal::namespace_short)1>>
 // CHECK: template <> struct KernelInfo<::dummy_functor_4<(enum_in_anonNS)1>>
 // CHECK: template <> struct KernelInfo<::dummy_functor_5<(no_type_set)0>>
+// CHECK: template <> struct KernelInfo<::dummy_functor_6<(unscoped_enum)0>>
+// CHECK: template <> struct KernelInfo<::dummy_functor_7<(unscoped_enum_no_type_set)1>>

From 46ef79d9d48283ce2a20350d4e21b391fd8b168a Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews@intel.com>
Date: Thu, 14 May 2020 15:58:04 -0700
Subject: [PATCH 07/14] Refactored printArguments

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>
---
 clang/lib/Sema/SemaSYCL.cpp | 91 +++++++++++++++++++++----------------
 1 file changed, 51 insertions(+), 40 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index fdb8940e31e15..a19432c3d32b3 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -1845,55 +1845,66 @@ static std::string getCPPTypeString(QualType Ty) {
 
 static void printArguments(ASTContext &Ctx, raw_ostream &ArgOS,
                            ArrayRef<TemplateArgument> Args,
-                           const PrintingPolicy &P, bool ParameterPack) {
+                           const PrintingPolicy &P);
 
-  if (!ParameterPack)
-    ArgOS << "<";
+static void printArgument(ASTContext &Ctx, raw_ostream &ArgOS,
+                          TemplateArgument Arg, const PrintingPolicy &P) {
+  switch (Arg.getKind()) {
+  case TemplateArgument::ArgKind::Pack: {
+    printArguments(Ctx, ArgOS, Arg.getPackAsArray(), P);
+    break;
+  }
+  case TemplateArgument::ArgKind::Integral: {
+    QualType T = Arg.getIntegralType();
+    const EnumType *ET = T->getAs<EnumType>();
+
+    if (ET) {
+      const llvm::APSInt &Val = Arg.getAsIntegral();
+      ArgOS << "(" << ET->getDecl()->getQualifiedNameAsString() << ")" << Val;
+    } else {
+      Arg.print(P, ArgOS);
+    }
+    break;
+  }
+  case TemplateArgument::ArgKind::Type: {
+    LangOptions LO;
+    PrintingPolicy TypePolicy(LO);
+    TypePolicy.SuppressTypedefs = true;
+    TypePolicy.SuppressTagKeyword = true;
+    QualType T = Arg.getAsType();
+    QualType FullyQualifiedType = TypeName::getFullyQualifiedType(T, Ctx, true);
+    ArgOS << FullyQualifiedType.getAsString(TypePolicy);
+    break;
+  }
+  default:
+    Arg.print(P, ArgOS);
+  }
+}
 
+static void printArguments(ASTContext &Ctx, raw_ostream &ArgOS,
+                           ArrayRef<TemplateArgument> Args,
+                           const PrintingPolicy &P) {
   bool FirstArg = true;
   const char *Comma = ", ";
 
   for (unsigned I = 0; I < Args.size(); I++) {
     const TemplateArgument &Arg = Args[I];
-    if (Arg.getKind() == TemplateArgument::ArgKind::Pack) {
-      if (Arg.pack_size() && !FirstArg)
-        ArgOS << Comma;
-      printArguments(Ctx, ArgOS, Arg.getPackAsArray(), P,
-                     /*ParameterPack*/ true);
-    } else if (Arg.getKind() == TemplateArgument::ArgKind::Integral) {
-      if (!FirstArg)
-        ArgOS << Comma;
-
-      QualType T = Arg.getIntegralType();
-      const EnumType *ET = T->getAs<EnumType>();
-
-      if (ET) {
-        const llvm::APSInt &Val = Arg.getAsIntegral();
-        ArgOS << "(" << ET->getDecl()->getQualifiedNameAsString() << ")" << Val;
-      } else {
-        Arg.print(P, ArgOS);
-      }
-    } else if (Arg.getKind() == TemplateArgument::ArgKind::Type) {
-      if (!FirstArg)
-        ArgOS << Comma;
-      LangOptions LO;
-      PrintingPolicy TypePolicy(LO);
-      TypePolicy.SuppressTypedefs = true;
-      TypePolicy.SuppressTagKeyword = true;
-      QualType T = Arg.getAsType();
-      QualType FullyQualifiedType =
-          TypeName::getFullyQualifiedType(T, Ctx, true);
-      ArgOS << FullyQualifiedType.getAsString(TypePolicy);
-    } else {
-      if (!FirstArg)
-        ArgOS << Comma;
-      Arg.print(P, ArgOS);
-    }
+
+    if (!FirstArg)
+      ArgOS << Comma;
+
+    printArgument(Ctx, ArgOS, Arg, P);
+
     FirstArg = false;
   }
+}
 
-  if (!ParameterPack)
-    ArgOS << ">";
+static void printTemplateArguments(ASTContext &Ctx, raw_ostream &ArgOS,
+                                   ArrayRef<TemplateArgument> Args,
+                                   const PrintingPolicy &P) {
+  ArgOS << "<";
+  printArguments(Ctx, ArgOS, Args, P);
+  ArgOS << ">";
 }
 
 static std::string getKernelNameTypeString(QualType T) {
@@ -1920,7 +1931,7 @@ static std::string getKernelNameTypeString(QualType T) {
     // Print template arguments substituting enumerators
     ASTContext &Ctx = RD->getASTContext();
     const TemplateArgumentList &Args = TSD->getTemplateArgs();
-    printArguments(Ctx, ArgOS, Args.asArray(), P, /*ParameterPack*/ false);
+    printTemplateArguments(Ctx, ArgOS, Args.asArray(), P);
 
     return eraseAnonNamespace(ArgOS.str().str());
   }

From 6262d69cd519dd982c0d4bcd028d4b594d9e5005 Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews@intel.com>
Date: Thu, 14 May 2020 16:26:56 -0700
Subject: [PATCH 08/14] Minor refactoring

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>
---
 clang/lib/Sema/SemaSYCL.cpp | 7 ++-----
 1 file changed, 2 insertions(+), 5 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index a19432c3d32b3..4446fcb775f9e 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -1884,14 +1884,11 @@ static void printArgument(ASTContext &Ctx, raw_ostream &ArgOS,
 static void printArguments(ASTContext &Ctx, raw_ostream &ArgOS,
                            ArrayRef<TemplateArgument> Args,
                            const PrintingPolicy &P) {
-  bool FirstArg = true;
-  const char *Comma = ", ";
-
   for (unsigned I = 0; I < Args.size(); I++) {
     const TemplateArgument &Arg = Args[I];
 
-    if (!FirstArg)
-      ArgOS << Comma;
+    if (I != 0)
+      ArgOS << ", ";
 
     printArgument(Ctx, ArgOS, Arg, P);
 

From b59e1a3ba263b39bf8edb444735ce649fa8169be Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews@intel.com>
Date: Thu, 14 May 2020 16:32:53 -0700
Subject: [PATCH 09/14] Missed a line in previous commit

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>
---
 clang/lib/Sema/SemaSYCL.cpp | 1 -
 1 file changed, 1 deletion(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 4446fcb775f9e..3f59395398528 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -1892,7 +1892,6 @@ static void printArguments(ASTContext &Ctx, raw_ostream &ArgOS,
 
     printArgument(Ctx, ArgOS, Arg, P);
 
-    FirstArg = false;
   }
 }
 

From 0a83e1691b184b9af6540c5526099835459f4ffc Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews@intel.com>
Date: Thu, 14 May 2020 16:38:57 -0700
Subject: [PATCH 10/14] Clang-Format Changes

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>
---
 clang/lib/Sema/SemaSYCL.cpp | 1 -
 1 file changed, 1 deletion(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 3f59395398528..14e22f39bc07a 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -1891,7 +1891,6 @@ static void printArguments(ASTContext &Ctx, raw_ostream &ArgOS,
       ArgOS << ", ";
 
     printArgument(Ctx, ArgOS, Arg, P);
-
   }
 }
 

From 52add49316f519ad91f8cea85ee48528443a39b6 Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews@intel.com>
Date: Fri, 15 May 2020 20:14:09 -0700
Subject: [PATCH 11/14] Added diagnostic to handle unscoped enums with no fixed
 underlying type.

Also fixed handling of template argument type - enum type. The prior
patch handled only enum values, not the type itself.

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>
---
 .../clang/Basic/DiagnosticSemaKinds.td        |  4 +-
 clang/lib/AST/Decl.cpp                        |  3 +-
 clang/lib/Sema/SemaSYCL.cpp                   | 35 +++++---
 clang/test/CodeGenSYCL/kernelname-enum.cpp    | 20 ++---
 clang/test/SemaSYCL/kernelname-enum.cpp       | 81 +++++++++++++++++++
 5 files changed, 120 insertions(+), 23 deletions(-)
 create mode 100644 clang/test/SemaSYCL/kernelname-enum.cpp

diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index f4f20dfe13966..9c8d06c5c9a17 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -10761,7 +10761,9 @@ def err_builtin_launder_invalid_arg : Error<
 // SYCL-specific diagnostics
 def err_sycl_kernel_incorrectly_named : Error<
   "kernel %select{name is missing"
-  "|needs to have a globally-visible name}0">;
+  "|needs to have a globally-visible name"
+  "|name cannot be templated using unscoped enum without fixed underlying type"
+  "}0">;
 def err_sycl_restrict : Error<
   "SYCL kernel cannot "
 "%select{use a non-const global variable"
diff --git a/clang/lib/AST/Decl.cpp b/clang/lib/AST/Decl.cpp
index e286388387a09..e4961c964ff47 100644
--- a/clang/lib/AST/Decl.cpp
+++ b/clang/lib/AST/Decl.cpp
@@ -1594,9 +1594,8 @@ void NamedDecl::printNestedNameSpecifier(raw_ostream &OS,
     Ctx = Ctx->getParent();
   }
 
-  if (WithGlobalNsPrefix) {
+  if (WithGlobalNsPrefix)
     OS << "::";
-  }
 
   for (const DeclContext *DC : llvm::reverse(Contexts)) {
     if (const auto *Spec = dyn_cast<ClassTemplateSpecializationDecl>(DC)) {
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 14e22f39bc07a..bca40291a7716 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -1637,6 +1637,18 @@ static std::string eraseAnonNamespace(std::string S) {
   return S;
 }
 
+static bool checkEnumTemplateParameter(const EnumDecl *ED,
+                                       DiagnosticsEngine &Diag,
+                                       SourceLocation KernelLocation) {
+  if (!ED->isScoped() && !ED->isFixed()) {
+    Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named) << 2;
+    Diag.Report(ED->getSourceRange().getBegin(), diag::note_entity_declared_at)
+        << ED;
+    return false;
+  }
+  return true;
+}
+
 // Emits a forward declaration
 void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
                                         SourceLocation KernelLocation) {
@@ -1774,8 +1786,20 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
 
       switch (Arg.getKind()) {
       case TemplateArgument::ArgKind::Type:
-        emitForwardClassDecls(O, Arg.getAsType(), KernelLocation, Printed);
+      case TemplateArgument::ArgKind::Integral: {
+        QualType T = (Arg.getKind() == TemplateArgument::ArgKind::Type)
+                         ? Arg.getAsType()
+                         : Arg.getIntegralType();
+
+        // Handle Kernel Name Type templated using enum type and value.
+        if (const auto *ET = T->getAs<EnumType>()) {
+          const EnumDecl *ED = ET->getDecl();
+          if (checkEnumTemplateParameter(ED, Diag, KernelLocation))
+            emitFwdDecl(O, ED, KernelLocation);
+        } else if (Arg.getKind() == TemplateArgument::ArgKind::Type)
+          emitForwardClassDecls(O, T, KernelLocation, Printed);
         break;
+      }
       case TemplateArgument::ArgKind::Pack: {
         ArrayRef<TemplateArgument> Pack = Arg.getPackAsArray();
 
@@ -1808,15 +1832,6 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
         }
         break;
       }
-      case TemplateArgument::ArgKind::Integral: {
-        // Handle Kernel Name Type templated using enum.
-        QualType T = Arg.getIntegralType();
-        if (const auto *ET = T->getAs<EnumType>()) {
-          const EnumDecl *ED = ET->getDecl();
-          emitFwdDecl(O, ED, KernelLocation);
-        }
-        break;
-      }
       default:
         break; // nop
       }
diff --git a/clang/test/CodeGenSYCL/kernelname-enum.cpp b/clang/test/CodeGenSYCL/kernelname-enum.cpp
index 8ae470a4f402b..8aa170b15a566 100644
--- a/clang/test/CodeGenSYCL/kernelname-enum.cpp
+++ b/clang/test/CodeGenSYCL/kernelname-enum.cpp
@@ -8,11 +8,6 @@ enum unscoped_enum : int {
   val_2
 };
 
-enum unscoped_enum_no_type_set {
-  val_3,
-  val_4
-};
-
 enum class no_namespace_int : int {
   val_1,
   val_2
@@ -78,7 +73,7 @@ class dummy_functor_6 {
   void operator()() {}
 };
 
-template <unscoped_enum_no_type_set EnumType>
+template <typename EnumType>
 class dummy_functor_7 {
 public:
   void operator()() {}
@@ -92,7 +87,8 @@ int main() {
   dummy_functor_4<enum_in_anonNS::val_2> f4;
   dummy_functor_5<no_type_set::val_1> f5;
   dummy_functor_6<unscoped_enum::val_1> f6;
-  dummy_functor_7<unscoped_enum_no_type_set::val_4> f7;
+  dummy_functor_7<no_namespace_int> f7;
+  dummy_functor_7<internal::namespace_short> f8;
 
   cl::sycl::queue q;
 
@@ -124,6 +120,10 @@ int main() {
     cgh.single_task(f7);
   });
 
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f8);
+  });
+
   return 0;
 }
 
@@ -144,8 +144,7 @@ int main() {
 // CHECK: template <no_type_set EnumType> class dummy_functor_5;
 // CHECK: enum unscoped_enum : int;
 // CHECK: template <unscoped_enum EnumType> class dummy_functor_6;
-// CHECK: enum unscoped_enum_no_type_set : unsigned int;
-// CHECK: template <unscoped_enum_no_type_set EnumType> class dummy_functor_7;
+// CHECK: template <typename EnumType> class dummy_functor_7;
 
 // CHECK: Specializations of KernelInfo for kernel function types:
 // CHECK: template <> struct KernelInfo<::dummy_functor_1<(no_namespace_int)0>>
@@ -154,4 +153,5 @@ int main() {
 // CHECK: template <> struct KernelInfo<::dummy_functor_4<(enum_in_anonNS)1>>
 // CHECK: template <> struct KernelInfo<::dummy_functor_5<(no_type_set)0>>
 // CHECK: template <> struct KernelInfo<::dummy_functor_6<(unscoped_enum)0>>
-// CHECK: template <> struct KernelInfo<::dummy_functor_7<(unscoped_enum_no_type_set)1>>
+// CHECK: template <> struct KernelInfo<::dummy_functor_7<::no_namespace_int>>
+// CHECK: template <> struct KernelInfo<::dummy_functor_7<::internal::namespace_short>>
diff --git a/clang/test/SemaSYCL/kernelname-enum.cpp b/clang/test/SemaSYCL/kernelname-enum.cpp
new file mode 100644
index 0000000000000..d3c7a982c8248
--- /dev/null
+++ b/clang/test/SemaSYCL/kernelname-enum.cpp
@@ -0,0 +1,81 @@
+// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s
+
+//#include <sycl.hpp>
+
+#include "sycl.hpp"
+
+enum unscoped_enum_int : int {
+  val_1,
+  val_2
+};
+
+// expected-note@+1 {{'unscoped_enum_no_type_set' declared here}}
+enum unscoped_enum_no_type_set {
+  val_3,
+  val_4
+};
+
+enum class scoped_enum_int : int {
+  val_1,
+  val_2
+};
+
+enum class scoped_enum_no_type_set {
+  val_3,
+  val_4
+};
+
+template <unscoped_enum_int EnumType>
+class dummy_functor_1 {
+public:
+  void operator()() {}
+};
+
+// expected-error@+2 {{kernel name cannot be templated using unscoped enum without fixed underlying type}}
+template <unscoped_enum_no_type_set EnumType>
+class dummy_functor_2 {
+public:
+  void operator()() {}
+};
+
+template <scoped_enum_int EnumType>
+class dummy_functor_3 {
+public:
+  void operator()() {}
+};
+
+template <scoped_enum_no_type_set EnumType>
+class dummy_functor_4 {
+public:
+  void operator()() {}
+};
+
+int main() {
+
+  dummy_functor_1<val_1> f1;
+  dummy_functor_2<val_3> f2;
+  dummy_functor_3<scoped_enum_int::val_2> f3;
+  dummy_functor_4<scoped_enum_no_type_set::val_4> f4;
+
+  cl::sycl::queue q;
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f1);
+  });
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f2);
+  });
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f3);
+  });
+
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f4);
+  });
+
+
+  return 0;
+}
+

From 9d070232d8cf7dd95dd77c73c0dd69d557c8048c Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews@intel.com>
Date: Fri, 15 May 2020 20:54:36 -0700
Subject: [PATCH 12/14] Clang Format changes

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>
---
 clang/test/SemaSYCL/kernelname-enum.cpp | 2 --
 1 file changed, 2 deletions(-)

diff --git a/clang/test/SemaSYCL/kernelname-enum.cpp b/clang/test/SemaSYCL/kernelname-enum.cpp
index d3c7a982c8248..dde692b0ce3ff 100644
--- a/clang/test/SemaSYCL/kernelname-enum.cpp
+++ b/clang/test/SemaSYCL/kernelname-enum.cpp
@@ -75,7 +75,5 @@ int main() {
     cgh.single_task(f4);
   });
 
-
   return 0;
 }
-

From eead6673e09b4a1926814733ba84806d2f30cb82 Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews@intel.com>
Date: Sun, 17 May 2020 07:49:37 -0700
Subject: [PATCH 13/14] Implemented review comments:

Changed 'check' function to return true on error.
Changed diagnostic message.

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>
---
 clang/include/clang/Basic/DiagnosticSemaKinds.td | 2 +-
 clang/lib/Sema/SemaSYCL.cpp                      | 6 +++---
 clang/test/SemaSYCL/kernelname-enum.cpp          | 2 +-
 3 files changed, 5 insertions(+), 5 deletions(-)

diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 9c8d06c5c9a17..9afff5381820d 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -10762,7 +10762,7 @@ def err_builtin_launder_invalid_arg : Error<
 def err_sycl_kernel_incorrectly_named : Error<
   "kernel %select{name is missing"
   "|needs to have a globally-visible name"
-  "|name cannot be templated using unscoped enum without fixed underlying type"
+  "|name is invalid. Unscoped enum requires fixed underlying type"
   "}0">;
 def err_sycl_restrict : Error<
   "SYCL kernel cannot "
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index bca40291a7716..3258278234ddc 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -1644,9 +1644,9 @@ static bool checkEnumTemplateParameter(const EnumDecl *ED,
     Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named) << 2;
     Diag.Report(ED->getSourceRange().getBegin(), diag::note_entity_declared_at)
         << ED;
-    return false;
+    return true;
   }
-  return true;
+  return false;
 }
 
 // Emits a forward declaration
@@ -1794,7 +1794,7 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
         // Handle Kernel Name Type templated using enum type and value.
         if (const auto *ET = T->getAs<EnumType>()) {
           const EnumDecl *ED = ET->getDecl();
-          if (checkEnumTemplateParameter(ED, Diag, KernelLocation))
+          if (!checkEnumTemplateParameter(ED, Diag, KernelLocation))
             emitFwdDecl(O, ED, KernelLocation);
         } else if (Arg.getKind() == TemplateArgument::ArgKind::Type)
           emitForwardClassDecls(O, T, KernelLocation, Printed);
diff --git a/clang/test/SemaSYCL/kernelname-enum.cpp b/clang/test/SemaSYCL/kernelname-enum.cpp
index dde692b0ce3ff..538ee892cb608 100644
--- a/clang/test/SemaSYCL/kernelname-enum.cpp
+++ b/clang/test/SemaSYCL/kernelname-enum.cpp
@@ -31,7 +31,7 @@ class dummy_functor_1 {
   void operator()() {}
 };
 
-// expected-error@+2 {{kernel name cannot be templated using unscoped enum without fixed underlying type}}
+// expected-error@+2 {{kernel name is invalid. Unscoped enum requires fixed underlying type}}
 template <unscoped_enum_no_type_set EnumType>
 class dummy_functor_2 {
 public:

From 1064c459bccea583910b98ebe04e4a84ab81f04a Mon Sep 17 00:00:00 2001
From: Elizabeth Andrews <elizabeth.andrews@intel.com>
Date: Mon, 18 May 2020 09:46:19 -0700
Subject: [PATCH 14/14] Removed commented code

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>
---
 clang/test/SemaSYCL/kernelname-enum.cpp | 2 --
 1 file changed, 2 deletions(-)

diff --git a/clang/test/SemaSYCL/kernelname-enum.cpp b/clang/test/SemaSYCL/kernelname-enum.cpp
index 538ee892cb608..22a9f96acc50a 100644
--- a/clang/test/SemaSYCL/kernelname-enum.cpp
+++ b/clang/test/SemaSYCL/kernelname-enum.cpp
@@ -1,7 +1,5 @@
 // RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s
 
-//#include <sycl.hpp>
-
 #include "sycl.hpp"
 
 enum unscoped_enum_int : int {