diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index b0459d97c01f0..358158b2fb31a 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/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([]() {}); - * declared within an anonymous namespace (at any level) - e.g., namespace foo { namespace { class Whatever; } } - h.single_task([]() {}); * declared within a function e.g., void foo() { struct S { int i; }; h.single_task([]() {}); } @@ -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(DeclCtx)) { const auto *NSDecl = cast(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(); } diff --git a/clang/test/CodeGenSYCL/int_header1.cpp b/clang/test/CodeGenSYCL/int_header1.cpp index 66a923a622d95..c97495bbe3631 100644 --- a/clang/test/CodeGenSYCL/int_header1.cpp +++ b/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 { // 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> { // CHECK:template <> struct KernelInfo<::nm1::KernelName9> { // CHECK:template <> struct KernelInfo<::nm1::KernelName3>> { @@ -51,6 +66,18 @@ namespace { template 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>( [=]() { 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>( + [=]() { acc.use(); }); + + // kernel name type is a class, declared in the anonymous namespace + kernel_single_task( + [=]() { acc.use(); }); + + // kernel name types declared in nested anonymous namespace + kernel_single_task( + [=]() { acc.use(); }); + + kernel_single_task( + [=]() { acc.use(); }); + // Kernel name type is a templated specialization class with empty template pack argument kernel_single_task>( [=]() { acc.use(); }); diff --git a/clang/test/CodeGenSYCL/kernelname-enum.cpp b/clang/test/CodeGenSYCL/kernelname-enum.cpp index f7641c5bf3691..a087edf629c0a 100644 --- a/clang/test/CodeGenSYCL/kernelname-enum.cpp +++ b/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 +class dummy_functor_4 { +public: + void operator()() const {} +}; + template class dummy_functor_5 { public: @@ -100,6 +113,7 @@ int main() { dummy_functor_1 f1; dummy_functor_2 f2; dummy_functor_3 f3; + dummy_functor_4 f4; dummy_functor_5 f5; dummy_functor_6 f6; dummy_functor_7 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 class dummy_functor_3; +// CHECK: namespace { +// CHECK-NEXT: enum class enum_in_anonNS : short; +// CHECK-NEXT: } +// CHECK: template class dummy_functor_4; // CHECK: enum class no_type_set : int; // CHECK: template class dummy_functor_5; // CHECK: enum unscoped_enum : int; @@ -178,6 +200,7 @@ int main() { // CHECK: template <> struct KernelInfo<::dummy_functor_1(0)>> // CHECK: template <> struct KernelInfo<::dummy_functor_2(1)>> // CHECK: template <> struct KernelInfo<::dummy_functor_3(1)>> +// CHECK: template <> struct KernelInfo<::dummy_functor_4(1)>> // CHECK: template <> struct KernelInfo<::dummy_functor_5(0)>> // CHECK: template <> struct KernelInfo<::dummy_functor_6(0)>> // CHECK: template <> struct KernelInfo<::dummy_functor_7<::no_namespace_int>> diff --git a/clang/test/SemaSYCL/nested-anon-and-std-ns.cpp b/clang/test/SemaSYCL/nested-anon-and-std-ns.cpp index 5b8c58507637e..e4aca4f184ca9 100644 --- a/clang/test/SemaSYCL/nested-anon-and-std-ns.cpp +++ b/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([] {}); }); - // 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([] {}); - }); - - // 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([] {}); - }); - // no error for valid ns q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); diff --git a/sycl/test/functor/kernel_functor.cpp b/sycl/test/functor/kernel_functor.cpp index a88cb96ed0e65..467ea5b2d3032 100644 --- a/sycl/test/functor/kernel_functor.cpp +++ b/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 &Acc_) + : X(X_), Acc(Acc_) {} + + void operator()() const { Acc[0] += X; } + +private: + int X; + cl::sycl::accessor 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 Buf(A, 1); + Q.submit([&](cl::sycl::handler &cgh) { + auto Acc = Buf.get_access(cgh); + Functor1 F(X, Acc); + + cgh.single_task(F); + }); + Q.submit([&](cl::sycl::handler &cgh) { auto Acc = Buf.get_access(cgh); ns::Functor2 F(X, Acc); @@ -141,7 +169,7 @@ template 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);