From 98a1d7d109e641c819e852a0240f9918dd87af26 Mon Sep 17 00:00:00 2001 From: "srividya.sundaram" Date: Wed, 28 Apr 2021 13:05:28 -0700 Subject: [PATCH 1/5] [SYCL] Allow kernel names in anonymous namespace --- clang/lib/Sema/SemaSYCL.cpp | 13 +------- clang/test/CodeGenSYCL/int_header1.cpp | 6 ++++ clang/test/CodeGenSYCL/kernelname-enum.cpp | 23 ++++++++++++++ .../test/SemaSYCL/nested-anon-and-std-ns.cpp | 12 -------- sycl/test/functor/kernel_functor.cpp | 30 ++++++++++++++++++- 5 files changed, 59 insertions(+), 25 deletions(-) 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..8eff8cfd016ca 100644 --- a/clang/test/CodeGenSYCL/int_header1.cpp +++ b/clang/test/CodeGenSYCL/int_header1.cpp @@ -9,6 +9,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>> { @@ -113,6 +114,11 @@ 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 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..395da62faa56f 100644 --- a/clang/test/SemaSYCL/nested-anon-and-std-ns.cpp +++ b/clang/test/SemaSYCL/nested-anon-and-std-ns.cpp @@ -10,12 +10,6 @@ struct NestedStruct {}; }; // namespace NestedInStd }; // namespace std -namespace { -namespace NestedInAnon { -struct StructInAnonymousNS {}; -} // namespace NestedInAnon -} // namespace - namespace Named { namespace { struct IsThisValid {}; @@ -50,12 +44,6 @@ struct MyWrapper { 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); From d75185f2700891d97b176bb94c6d6ff75fdc8afd Mon Sep 17 00:00:00 2001 From: "srividya.sundaram" Date: Wed, 28 Apr 2021 14:27:56 -0700 Subject: [PATCH 2/5] update test --- clang/test/SemaSYCL/nested-anon-and-std-ns.cpp | 12 ------------ 1 file changed, 12 deletions(-) diff --git a/clang/test/SemaSYCL/nested-anon-and-std-ns.cpp b/clang/test/SemaSYCL/nested-anon-and-std-ns.cpp index 395da62faa56f..e4aca4f184ca9 100644 --- a/clang/test/SemaSYCL/nested-anon-and-std-ns.cpp +++ b/clang/test/SemaSYCL/nested-anon-and-std-ns.cpp @@ -10,12 +10,6 @@ struct NestedStruct {}; }; // namespace NestedInStd }; // namespace std -namespace Named { -namespace { -struct IsThisValid {}; -} // namespace -} // namespace Named - namespace ValidNS { struct StructinValidNS {}; } // namespace ValidNS @@ -38,12 +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([] {}); - }); - // no error for valid ns q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); From 4b26c05efd9e59f47577350cf12a5b53e5e779fb Mon Sep 17 00:00:00 2001 From: "srividya.sundaram" Date: Thu, 29 Apr 2021 15:05:52 -0700 Subject: [PATCH 3/5] Update int_header1.cpp test --- clang/test/CodeGenSYCL/int_header1.cpp | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/clang/test/CodeGenSYCL/int_header1.cpp b/clang/test/CodeGenSYCL/int_header1.cpp index 8eff8cfd016ca..6c496941b8be3 100644 --- a/clang/test/CodeGenSYCL/int_header1.cpp +++ b/clang/test/CodeGenSYCL/int_header1.cpp @@ -1,6 +1,12 @@ // 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:class ClassInAnonNS; +// CHECK:} + // CHECK:template <> struct KernelInfo { // CHECK:template <> struct KernelInfo<::nm1::nm2::KernelName0> { // CHECK:template <> struct KernelInfo<::nm1::KernelName1> { @@ -119,6 +125,10 @@ struct MyWrapper { kernel_single_task>( [=]() { acc.use(); }); + // kernel name type is a class, declared in the anonymous namespace + kernel_single_task( + [=]() { acc.use(); }); + // Kernel name type is a templated specialization class with empty template pack argument kernel_single_task>( [=]() { acc.use(); }); From c19761420b362a4b5acdbe1f4bb6a10a37f510ad Mon Sep 17 00:00:00 2001 From: "srividya.sundaram" Date: Thu, 29 Apr 2021 22:05:14 -0700 Subject: [PATCH 4/5] Add check-next to int_header1.cpp --- clang/test/CodeGenSYCL/int_header1.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/test/CodeGenSYCL/int_header1.cpp b/clang/test/CodeGenSYCL/int_header1.cpp index 6c496941b8be3..14a36fb03510d 100644 --- a/clang/test/CodeGenSYCL/int_header1.cpp +++ b/clang/test/CodeGenSYCL/int_header1.cpp @@ -4,8 +4,8 @@ // Check if forward declarations of kernel names in anonymous namespace are in // anonymous namespace in the integration header as well. // CHECK:namespace { -// CHECK:class ClassInAnonNS; -// CHECK:} +// CHECK-NEXT:class ClassInAnonNS; +// CHECK-NEXT:} // CHECK:template <> struct KernelInfo { // CHECK:template <> struct KernelInfo<::nm1::nm2::KernelName0> { From 705af36cde7fccb131eef5e3ce522bd9de43f83c Mon Sep 17 00:00:00 2001 From: "srividya.sundaram" Date: Mon, 3 May 2021 09:56:30 -0700 Subject: [PATCH 5/5] Update test with nested anon ns cases --- clang/test/CodeGenSYCL/int_header1.cpp | 27 ++++++++++++++++++++++++++ 1 file changed, 27 insertions(+) diff --git a/clang/test/CodeGenSYCL/int_header1.cpp b/clang/test/CodeGenSYCL/int_header1.cpp index 14a36fb03510d..c97495bbe3631 100644 --- a/clang/test/CodeGenSYCL/int_header1.cpp +++ b/clang/test/CodeGenSYCL/int_header1.cpp @@ -7,6 +7,14 @@ // 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> { @@ -58,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 {}; @@ -129,6 +149,13 @@ struct MyWrapper { 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(); });