diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 1909c1d81a872..2ab770ca1f40a 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -3025,19 +3025,34 @@ class SYCLKernelNameTypeVisitor S.getASTContext().getLangOpts().SYCLUnnamedLambda; const DeclContext *DeclCtx = Tag->getDeclContext(); if (DeclCtx && !UnnamedLambdaEnabled) { - auto *NameSpace = dyn_cast_or_null(DeclCtx); - if (NameSpace && NameSpace->isStdNamespace()) { - S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named) - << KernelNameType; - S.Diag(KernelInvocationFuncLoc, diag::note_invalid_type_in_sycl_kernel) - << /* kernel name cannot be a type in the std namespace */ 2 - << QualType(Tag->getTypeForDecl(), 0); - IsInvalid = true; - return; - } - if (!DeclCtx->isTranslationUnit() && !isa(DeclCtx)) { - const bool KernelNameIsMissing = Tag->getName().empty(); - if (KernelNameIsMissing) { + + while (!DeclCtx->isTranslationUnit()) { + auto *NSDecl = dyn_cast_or_null(DeclCtx); + if (NSDecl && NSDecl->isStdNamespace()) { + S.Diag(KernelInvocationFuncLoc, + diag::err_sycl_kernel_incorrectly_named) + << KernelNameType; + S.Diag(KernelInvocationFuncLoc, + diag::note_invalid_type_in_sycl_kernel) + << /* kernel name cannot be a type in the std namespace */ 2 + << QualType(Tag->getTypeForDecl(), 0); + IsInvalid = true; + return; + } + if (NSDecl && NSDecl->isAnonymousNamespace()) { + S.Diag(KernelInvocationFuncLoc, + diag::err_sycl_kernel_incorrectly_named) + << KernelNameType; + S.Diag(KernelInvocationFuncLoc, + diag::note_invalid_type_in_sycl_kernel) + << /* kernel name is not globally-visible */ 0 + << QualType(Tag->getTypeForDecl(), 0); + IsInvalid = true; + return; + } + + const bool UnnamedTypeUsed = Tag->getName().empty(); + if (UnnamedTypeUsed) { S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named) << KernelNameType; @@ -3047,7 +3062,7 @@ class SYCLKernelNameTypeVisitor IsInvalid = true; return; } - if (Tag->isCompleteDefinition()) { + if (isa(DeclCtx) && Tag->isCompleteDefinition()) { S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named) << KernelNameType; @@ -3056,11 +3071,16 @@ class SYCLKernelNameTypeVisitor << /* kernel name is not globally-visible */ 0 << QualType(Tag->getTypeForDecl(), 0); IsInvalid = true; - } else { + return; + } + if (isa(DeclCtx) && !Tag->isCompleteDefinition()) { S.Diag(KernelInvocationFuncLoc, diag::warn_sycl_implicit_decl); S.Diag(Tag->getSourceRange().getBegin(), diag::note_previous_decl) << Tag->getName(); } + + // Repeat the above checks for DeclCtx's in the parent-chain + DeclCtx = DeclCtx->getParent(); } } } diff --git a/clang/test/CodeGenSYCL/int_header1.cpp b/clang/test/CodeGenSYCL/int_header1.cpp index be932f13820c1..b72d5022e4fe6 100644 --- a/clang/test/CodeGenSYCL/int_header1.cpp +++ b/clang/test/CodeGenSYCL/int_header1.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h %s -o %t.out +// RUN: %clang_cc1 -fsycl -fsycl-is-device -sycl-std=2020 -fsycl-int-header=%t.h %s -o %t.out // RUN: FileCheck -input-file=%t.h %s // CHECK:template <> struct KernelInfo { @@ -11,7 +11,6 @@ // CHECK:template <> struct KernelInfo<::nm1::KernelName3> { // CHECK:template <> struct KernelInfo<::nm1::KernelName4> { // 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>> { @@ -49,11 +48,6 @@ namespace nm1 { } // namespace nm1 -namespace { - class ClassInAnonNS; - template class TmplClassInAnonNS; -} - struct MyWrapper { class KN101 {}; @@ -130,8 +124,6 @@ struct MyWrapper { // 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>( @@ -165,7 +157,7 @@ int main() { KernelInfo>::getName(); KernelInfo>::getName(); KernelInfo>::getName(); - KernelInfo>::getName(); + KernelInfo>::getName(); #endif //__SYCL_DEVICE_ONLY__ } diff --git a/clang/test/SemaSYCL/allow-nested-structs.cpp b/clang/test/SemaSYCL/allow-nested-structs.cpp new file mode 100644 index 0000000000000..f2fc0408c59fe --- /dev/null +++ b/clang/test/SemaSYCL/allow-nested-structs.cpp @@ -0,0 +1,69 @@ +// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2020 -verify %s + +#include "sycl.hpp" + +struct NestedStruct1 { + struct NestedStruct2 { + struct NestedStruct3 {}; + }; +}; + +namespace { +struct StructInAnonymousNS {}; +} // namespace + +namespace ValidNS { +struct StructinValidNS {}; +} // namespace ValidNS + +struct Parent { + using A = struct { + struct Child1 { + struct Child2 {}; + }; + }; +}; + +struct MyWrapper { + +public: + void test() { + cl::sycl::queue q; + struct StructInsideFunc {}; + + // no error + q.submit([&](cl::sycl::handler &h) { + h.single_task([] {}); + }); + + // no error + q.submit([&](cl::sycl::handler &h) { + h.single_task([] {}); + }); + + // no error + q.submit([&](cl::sycl::handler &h) { + h.single_task([] {}); + }); + + // expected-error@Inputs/sycl.hpp:220 {{'(anonymous namespace)::StructInAnonymousNS' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{'(anonymous namespace)::StructInAnonymousNS' should be globally-visible}} + // expected-note@+2{{in instantiation of function template specialization}} + q.submit([&](cl::sycl::handler &h) { + h.single_task([] {}); + }); + + // expected-error@Inputs/sycl.hpp:220 {{'StructInsideFunc' is an invalid kernel name type}} + // expected-note@Inputs/sycl.hpp:220 {{'StructInsideFunc' should be globally-visible}} + // expected-note@+2{{in instantiation of function template specialization}} + q.submit([&](cl::sycl::handler &h) { + h.single_task([] {}); + }); + } +}; + +int main() { + cl::sycl::queue q; + + return 0; +} diff --git a/clang/test/SemaSYCL/unnamed-kernel.cpp b/clang/test/SemaSYCL/unnamed-kernel.cpp index c4e3746f78679..c0d03dd415107 100644 --- a/clang/test/SemaSYCL/unnamed-kernel.cpp +++ b/clang/test/SemaSYCL/unnamed-kernel.cpp @@ -52,20 +52,12 @@ struct MyWrapper { h.single_task>([] {}); }); -#ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName0' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName0' should be globally-visible}} - // expected-note@+3{{in instantiation of function template specialization}} -#endif + // no error q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); }); -#ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:220 {{'namespace1::KernelName' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName3' should be globally-visible}} - // expected-note@+3{{in instantiation of function template specialization}} -#endif + //no error q.submit([&](cl::sycl::handler &h) { h.single_task>([] {}); }); @@ -85,21 +77,13 @@ struct MyWrapper { }); using InvalidAlias = InvalidKernelName4; -#ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName4' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName4' should be globally-visible}} - // expected-note@+3{{in instantiation of function template specialization}} -#endif + // no error q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); }); using InvalidAlias1 = InvalidKernelName5; -#ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@Inputs/sycl.hpp:220 {{'namespace1::KernelName' is an invalid kernel name type}} - // expected-note@Inputs/sycl.hpp:220 {{'MyWrapper::InvalidKernelName5' should be globally-visible}} - // expected-note@+3{{in instantiation of function template specialization}} -#endif + // no error q.submit([&](cl::sycl::handler &h) { h.single_task>([] {}); }); @@ -124,4 +108,4 @@ int main() { q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); }); return 0; -} +} \ No newline at end of file diff --git a/sycl/test/functor/kernel_functor.cpp b/sycl/test/functor/kernel_functor.cpp index fe70853226696..6a9866ad19e5c 100644 --- a/sycl/test/functor/kernel_functor.cpp +++ b/sycl/test/functor/kernel_functor.cpp @@ -18,27 +18,6 @@ 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; -}; -} - -// Case 2: // - functor class is defined in a namespace // - the '()' operator: // * does not have parameters (to be used in 'single_task'). @@ -60,7 +39,7 @@ class Functor2 { }; } -// Case 3: +// Case 2: // - functor class is templated and defined in the translation unit scope // - the '()' operator: // * has a parameter of type cl::sycl::id<1> (to be used in 'parallel_for'). @@ -78,7 +57,7 @@ template class TmplFunctor { cl::sycl::accessor Acc; }; -// Case 4: +// Case 3: // - functor class is templated and defined in the translation unit scope // - the '()' operator: // * has a parameter of type cl::sycl::id<1> (to be used in 'parallel_for'). @@ -103,12 +82,6 @@ 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); @@ -167,7 +140,7 @@ template T bar(T X) { int main() { const int Res1 = foo(10); const int Res2 = bar(10); - const int Gold1 = 40; + const int Gold1 = 30; const int Gold2 = 80; assert(Res1 == Gold1);