[SYCL] Allow kernel names in anonymous namespace

clang/lib/Sema/SemaSYCL.cpp

@@ -3272,9 +3272,6 @@ class SYCLKernelNameTypeVisitor
       * declared within namespace 'std' (at any level)
         e.g., namespace std { namespace literals { class Whatever; } }
         h.single_task<std::literals::Whatever>([]() {});
-      * declared within an anonymous namespace (at any level)
-        e.g., namespace foo { namespace { class Whatever; } }
-        h.single_task<foo::Whatever>([]() {});
       * declared within a function
         e.g., void foo() { struct S { int i; };
         h.single_task<S>([]() {}); }
@@ -3298,7 +3295,7 @@ class SYCLKernelNameTypeVisitor
     if (DeclCtx && !UnnamedLambdaEnabled) {
 
       // Check if the kernel name declaration is declared within namespace
-      // "std" or "anonymous" namespace (at any level).
+      // "std" (at any level).
       while (!DeclCtx->isTranslationUnit() && isa<NamespaceDecl>(DeclCtx)) {
         const auto *NSDecl = cast<NamespaceDecl>(DeclCtx);
         if (NSDecl->isStdNamespace()) {
@@ -3308,14 +3305,6 @@ class SYCLKernelNameTypeVisitor
           IsInvalid = true;
           return;
         }
-        if (NSDecl->isAnonymousNamespace()) {
-          S.Diag(KernelInvocationFuncLoc,
-                 diag::err_sycl_kernel_incorrectly_named)
-              << /* kernel name should be globally visible */ 0
-              << KernelNameType;
-          IsInvalid = true;
-          return;
-        }
         DeclCtx = DeclCtx->getParent();
       }
 

clang/test/CodeGenSYCL/int_header1.cpp

@@ -1,6 +1,20 @@
 // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -fsycl-int-header=%t.h %s -o %t.out
 // RUN: FileCheck -input-file=%t.h %s
 
+// Check if forward declarations of kernel names in anonymous namespace are in
+// anonymous namespace in the integration header as well.
+// CHECK:namespace  {
+// CHECK-NEXT:class ClassInAnonNS;
+// CHECK-NEXT:}
+
+// CHECK:namespace  { namespace NestedInAnon {
+// CHECK-NEXT:struct StructInAnonymousNS;
+// CHECK-NEXT:}}
+
+// CHECK:namespace Named { namespace  {
+// CHECK-NEXT:struct IsThisValid;
+// CHECK-NEXT:}}
+
 // CHECK:template <> struct KernelInfo<class KernelName> {
 // CHECK:template <> struct KernelInfo<::nm1::nm2::KernelName0> {
 // CHECK:template <> struct KernelInfo<::nm1::KernelName1> {
@@ -9,6 +23,7 @@
 // CHECK:template <> struct KernelInfo<::nm1::KernelName4<::nm1::nm2::KernelName0>> {
 // CHECK:template <> struct KernelInfo<::nm1::KernelName4<::nm1::KernelName1>> {
 // CHECK:template <> struct KernelInfo<::nm1::KernelName8<::nm1::nm2::C>> {
+// CHECK:template <> struct KernelInfo<::TmplClassInAnonNS<ClassInAnonNS>> {
 // CHECK:template <> struct KernelInfo<::nm1::KernelName9<char>> {
 // CHECK:template <> struct KernelInfo<::nm1::KernelName3<const volatile ::nm1::KernelName3<const volatile char>>> {
 
@@ -51,6 +66,18 @@ namespace {
   template <typename T> class TmplClassInAnonNS;
 }
 
+namespace {
+namespace NestedInAnon {
+struct StructInAnonymousNS {};
+} // namespace NestedInAnon
+} // namespace
+
+namespace Named {
+namespace {
+struct IsThisValid {};
+} // namespace
+} // namespace Named
+
 struct MyWrapper {
   class KN101 {};
 
@@ -113,6 +140,22 @@ struct MyWrapper {
     kernel_single_task<nm1::KernelName8<nm1::nm2::C>>(
       [=]() { acc.use(); });
 
+    // kernel name type is a templated class, both the top-level class and the
+    // template argument are declared in the anonymous namespace
+    kernel_single_task<TmplClassInAnonNS<class ClassInAnonNS>>(
+        [=]() { acc.use(); });
+
+    // kernel name type is a class, declared in the anonymous namespace
+    kernel_single_task<ClassInAnonNS>(
+        [=]() { acc.use(); });
+
+    // kernel name types declared in nested anonymous namespace
+    kernel_single_task<NestedInAnon::StructInAnonymousNS>(
+        [=]() { acc.use(); });
+
+    kernel_single_task<Named::IsThisValid>(
+        [=]() { acc.use(); });
+
     // Kernel name type is a templated specialization class with empty template pack argument
     kernel_single_task<nm1::KernelName9<char>>(
         [=]() { acc.use(); });

clang/test/CodeGenSYCL/kernelname-enum.cpp

@@ -25,6 +25,13 @@ enum class namespace_short : short {
 };
 }
 
+namespace {
+enum class enum_in_anonNS : short {
+  val_1,
+  val_2
+};
+}
+
 enum class no_type_set {
   val_1,
   val_2
@@ -48,6 +55,12 @@ class dummy_functor_3 {
   void operator()() const {}
 };
 
+template <enum_in_anonNS EnumType>
+class dummy_functor_4 {
+public:
+  void operator()() const {}
+};
+
 template <no_type_set EnumType>
 class dummy_functor_5 {
 public:
@@ -100,6 +113,7 @@ 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;
   dummy_functor_6<unscoped_enum::val_1> f6;
   dummy_functor_7<no_namespace_int> f7;
@@ -120,6 +134,10 @@ int main() {
     cgh.single_task(f3);
   });
 
+  q.submit([&](cl::sycl::handler &cgh) {
+    cgh.single_task(f4);
+  });
+
   q.submit([&](cl::sycl::handler &cgh) {
     cgh.single_task(f5);
   });
@@ -160,6 +178,10 @@ 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: enum class no_type_set : int;
 // CHECK: template <no_type_set EnumType> class dummy_functor_5;
 // CHECK: enum unscoped_enum : int;
@@ -178,6 +200,7 @@ int main() {
 // CHECK: template <> struct KernelInfo<::dummy_functor_1<static_cast<::no_namespace_int>(0)>>
 // CHECK: template <> struct KernelInfo<::dummy_functor_2<static_cast<::no_namespace_short>(1)>>
 // CHECK: template <> struct KernelInfo<::dummy_functor_3<static_cast<::internal::namespace_short>(1)>>
+// CHECK: template <> struct KernelInfo<::dummy_functor_4<static_cast<::enum_in_anonNS>(1)>>
 // CHECK: template <> struct KernelInfo<::dummy_functor_5<static_cast<::no_type_set>(0)>>
 // CHECK: template <> struct KernelInfo<::dummy_functor_6<static_cast<::unscoped_enum>(0)>>
 // CHECK: template <> struct KernelInfo<::dummy_functor_7<::no_namespace_int>>

clang/test/SemaSYCL/nested-anon-and-std-ns.cpp

@@ -10,18 +10,6 @@ struct NestedStruct {};
 }; // namespace NestedInStd
 }; // namespace std
 
-namespace {
-namespace NestedInAnon {
-struct StructInAnonymousNS {};
-} // namespace NestedInAnon
-} // namespace
-
-namespace Named {
-namespace {
-struct IsThisValid {};
-} // namespace
-} // namespace Named
-
 namespace ValidNS {
 struct StructinValidNS {};
 } // namespace ValidNS
@@ -44,18 +32,6 @@ struct MyWrapper {
       h.single_task<std::NestedInStd::NestedStruct>([] {});
     });
 
-    // expected-error@#KernelSingleTask {{'Named::(anonymous namespace)::IsThisValid' should be globally visible}}
-    // expected-note@+2{{in instantiation of function template specialization}}
-    q.submit([&](cl::sycl::handler &h) {
-      h.single_task<Named::IsThisValid>([] {});
-    });
-
-    // expected-error@#KernelSingleTask {{'(anonymous namespace)::NestedInAnon::StructInAnonymousNS' should be globally visible}}
-    // expected-note@+2{{in instantiation of function template specialization}}
-    q.submit([&](cl::sycl::handler &h) {
-      h.single_task<NestedInAnon::StructInAnonymousNS>([] {});
-    });
-
     // no error for valid ns
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<ValidNS::StructinValidNS>([] {});

sycl/test/functor/kernel_functor.cpp

@@ -18,6 +18,27 @@
 constexpr auto sycl_read_write = cl::sycl::access::mode::read_write;
 constexpr auto sycl_global_buffer = cl::sycl::access::target::global_buffer;
 
+// Case 1:
+// - functor class is defined in an anonymous namespace
+// - the '()' operator:
+//   * does not have parameters (to be used in 'single_task').
+//   * has the 'const' qualifier
+namespace {
+class Functor1 {
+public:
+  Functor1(
+      int X_,
+      cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> &Acc_)
+      : X(X_), Acc(Acc_) {}
+
+  void operator()() const { Acc[0] += X; }
+
+private:
+  int X;
+  cl::sycl::accessor<int, 1, sycl_read_write, sycl_global_buffer> Acc;
+};
+} // namespace
+
 // Case 1:
 // - functor class is defined in a namespace
 // - the '()' operator:
@@ -83,6 +104,13 @@ int foo(int X) {
     cl::sycl::queue Q;
     cl::sycl::buffer<int, 1> Buf(A, 1);
 
+    Q.submit([&](cl::sycl::handler &cgh) {
+      auto Acc = Buf.get_access<sycl_read_write, sycl_global_buffer>(cgh);
+      Functor1 F(X, Acc);
+
+      cgh.single_task(F);
+    });
+
     Q.submit([&](cl::sycl::handler &cgh) {
       auto Acc = Buf.get_access<sycl_read_write, sycl_global_buffer>(cgh);
       ns::Functor2 F(X, Acc);
@@ -141,7 +169,7 @@ template <typename T> T bar(T X) {
 int main() {
   const int Res1 = foo(10);
   const int Res2 = bar(10);
-  const int Gold1 = 30;
+  const int Gold1 = 40;
   const int Gold2 = 80;
 
   assert(Res1 == Gold1);