[SYCL] Add support for kernel name types templated using enums.

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"

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

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
       }

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

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