diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index d87176ec939c8..32a8407cbe333 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -7219,9 +7219,11 @@ def warn_format_nonliteral : Warning< InGroup, DefaultIgnore; def err_sycl_device_global_incorrect_scope : Error< - "'device_global' variables must be static or declared at namespace scope">; + "'device_global' variable must be a static data member or declared in global or namespace scope">; def err_sycl_device_global_not_publicly_accessible: Error< - "'device_global' member variable %0 is not publicly accessible from namespace scope">; + "'device_global' member variable %0 should be publicly accessible from namespace scope">; +def err_sycl_device_global_array : Error< + "'device_global' array is not allowed">; def err_unexpected_interface : Error< "unexpected interface name %0: expected expression">; diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index f337d0f2df173..cbc56e8cd2316 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -7678,13 +7678,32 @@ NamedDecl *Sema::ActOnVariableDeclarator( NewVD->setTSCSpec(TSCS); } - // Global variables with types decorated with device_global attribute must be - // static if they are declared in SYCL device code. if (getLangOpts().SYCLIsDevice) { - if (SCSpec != DeclSpec::SCS_static && !NewVD->hasGlobalStorage() && - isTypeDecoratedWithDeclAttribute( - NewVD->getType())) - Diag(D.getIdentifierLoc(), diag::err_sycl_device_global_incorrect_scope); + // device_global array is not allowed. + if (const ArrayType *AT = getASTContext().getAsArrayType(NewVD->getType())) + if (isTypeDecoratedWithDeclAttribute( + AT->getElementType())) + Diag(NewVD->getLocation(), diag::err_sycl_device_global_array); + + // Global variables with types decorated with device_global attribute must + // be static if they are declared in SYCL device code. + if (isTypeDecoratedWithDeclAttribute( + NewVD->getType())) { + if (SCSpec == DeclSpec::SCS_static) { + const DeclContext *DC = NewVD->getDeclContext(); + while (!DC->isTranslationUnit()) { + if (isa(DC)) { + Diag(D.getIdentifierLoc(), + diag::err_sycl_device_global_incorrect_scope); + break; + } + DC = DC->getParent(); + } + } else if (!NewVD->hasGlobalStorage()) { + Diag(D.getIdentifierLoc(), + diag::err_sycl_device_global_incorrect_scope); + } + } // Static variables declared inside SYCL device code must be const or // constexpr unless their types are decorated with global_variable_allowed diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp index 47976032f0a75..52aba2a9f4968 100644 --- a/clang/lib/Sema/SemaDeclCXX.cpp +++ b/clang/lib/Sema/SemaDeclCXX.cpp @@ -3598,10 +3598,24 @@ Sema::ActOnCXXMemberDeclarator(Scope *S, AccessSpecifier AS, Declarator &D, if (getLangOpts().SYCLIsDevice) { if (auto Value = dyn_cast(Member)) { if (isTypeDecoratedWithDeclAttribute( - Value->getType()) && - Value->getAccess() != AS_public) { - Diag(Loc, diag::err_sycl_device_global_not_publicly_accessible) - << Value; + Value->getType())) { + if (Value->getAccess() == AS_private || + Value->getAccess() == AS_protected) { + Diag(Loc, diag::err_sycl_device_global_not_publicly_accessible) + << Value; + } + const DeclContext *DC = Member->getDeclContext(); + while (!DC->isTranslationUnit()) { + if (auto Decl = dyn_cast(DC)) { + if (Decl->getAccess() == AS_private || + Decl->getAccess() == AS_protected) { + Diag(Loc, diag::err_sycl_device_global_not_publicly_accessible) + << Value; + break; + } + } + DC = DC->getParent(); + } } } } diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index 6b2a47f8ddd4d..efa8b26c5d9cc 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -1614,7 +1614,30 @@ Decl *TemplateDeclInstantiator::VisitVarDecl(VarDecl *D, // Only add this if we aren't instantiating a variable template. We'll end up // adding the VarTemplateSpecializationDecl later. if (!InstantiatingVarTemplate) { - SemaRef.addSyclVarDecl(Var); + if (SemaRef.getLangOpts().SYCLIsDevice && + SemaRef.isTypeDecoratedWithDeclAttribute( + Var->getType())) { + if (!Var->hasGlobalStorage()) + SemaRef.Diag(D->getLocation(), + diag::err_sycl_device_global_incorrect_scope); + + if (Var->getAccess() == AS_private || Var->getAccess() == AS_protected) + SemaRef.Diag(D->getLocation(), + diag::err_sycl_device_global_not_publicly_accessible) + << Var; + + if (Var->isStaticLocal()) { + const DeclContext *DC = Var->getDeclContext(); + while (!DC->isTranslationUnit()) { + if (isa(DC)) { + SemaRef.Diag(D->getLocation(), + diag::err_sycl_device_global_incorrect_scope); + break; + } + DC = DC->getParent(); + } + } + } if (const auto *SYCLDevice = Var->getAttr()) { if (!SemaRef.isTypeDecoratedWithDeclAttribute( Var->getType())) @@ -1622,6 +1645,7 @@ Decl *TemplateDeclInstantiator::VisitVarDecl(VarDecl *D, diag::err_sycl_attribute_not_device_global) << SYCLDevice; } + SemaRef.addSyclVarDecl(Var); } return Var; } @@ -1711,6 +1735,17 @@ Decl *TemplateDeclInstantiator::VisitFieldDecl(FieldDecl *D) { Field->setImplicit(D->isImplicit()); Field->setAccess(D->getAccess()); + // Static members are not processed here, so error out if we have a device + // global without checking access modifier. + if (SemaRef.getLangOpts().SYCLIsDevice) { + if (SemaRef.isTypeDecoratedWithDeclAttribute( + Field->getType())) { + SemaRef.Diag(D->getLocation(), + diag::err_sycl_device_global_incorrect_scope); + Field->setInvalidDecl(); + return nullptr; + } + } Owner->addDecl(Field); return Field; diff --git a/clang/test/SemaSYCL/device_global.cpp b/clang/test/SemaSYCL/device_global.cpp index 1a03a6c44c3ef..d4a4329f13667 100644 --- a/clang/test/SemaSYCL/device_global.cpp +++ b/clang/test/SemaSYCL/device_global.cpp @@ -1,9 +1,10 @@ // RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -verify %s #include "Inputs/sycl.hpp" -// Test cases below check for valid usage of device_global and -// global_variable_allowed attributes, and that they are being correctly -// generated in the AST. +// Diagnostic tests for device_global and global_variable_allowed attribute. + +// Test that there are no errors when variables of type device_global are +// decorated with global_variable_allowed attribute appropriately. using namespace sycl::ext::oneapi; device_global glob; // OK @@ -18,19 +19,95 @@ device_global Foo::d; struct Baz { private: - // expected-error@+1{{'device_global' member variable 'f' is not publicly accessible from namespace scope}} + // expected-error@+1{{'device_global' member variable 'f' should be publicly accessible from namespace scope}} static device_global f; + +protected: + // expected-error@+1{{'device_global' member variable 'g' should be publicly accessible from namespace scope}} + static device_global g; }; + device_global Baz::f; device_global not_array; // OK +// expected-error@+1{{'device_global' array is not allowed}} +device_global array[4]; + device_global same_name; // OK + namespace foo { device_global same_name; // OK } -namespace { -device_global same_name; // OK + +struct BBar { +private: + struct BarInsider { + // expected-error@+1{{'device_global' member variable 'c' should be publicly accessible from namespace scope}} + static device_global c; + }; + +protected: + struct BarInsiderProtected { + // expected-error@+1{{'device_global' member variable 'c' should be publicly accessible from namespace scope}} + static device_global c; + }; +}; + +struct ABar { + void method() { + // expected-error@+1{{'device_global' variable must be a static data member or declared in global or namespace scope}} + static device_global c; + } + struct BarInsider { + static device_global c; + void method() { + // expected-error@+1{{'device_global' variable must be a static data member or declared in global or namespace scope}} + static device_global c; + } + }; +}; + +template void fooBar() { + // expected-error@+1{{'device_global' variable must be a static data member or declared in global or namespace scope}} + static device_global c; + // expected-error@+1{{'device_global' variable must be a static data member or declared in global or namespace scope}} + device_global d; +} + +template struct TS { +private: + // expected-error@+1 2{{'device_global' member variable 'a' should be publicly accessible from namespace scope}} + static device_global a; + // expected-error@+1 2{{'device_global' variable must be a static data member or declared in global or namespace scope}} + device_global b; + // expected-error@+2{{'device_global' member variable 'c' should be publicly accessible from namespace scope}} + // expected-error@+1 2{{'device_global' variable must be a static data member or declared in global or namespace scope}} + device_global c; + +public: + static device_global d; + // expected-error@+1 2{{'device_global' variable must be a static data member or declared in global or namespace scope}} + device_global e; + // expected-error@+1 2{{'device_global' variable must be a static data member or declared in global or namespace scope}} + device_global f; + +protected: + // expected-error@+1 2{{'device_global' member variable 'g' should be publicly accessible from namespace scope}} + static device_global g; + // expected-error@+1 2{{'device_global' variable must be a static data member or declared in global or namespace scope}} + device_global h; + // expected-error@+2{{'device_global' member variable 'i' should be publicly accessible from namespace scope}} + // expected-error@+1 2{{'device_global' variable must be a static data member or declared in global or namespace scope}} + device_global i; +}; + +// expected-note@+1{{in instantiation of template class 'TS' requested here}} +TS AAAA; + +//expected-note@+2{{in instantiation of template class 'TS' requested here}} +template void templFoo () { + TS Var; } // expected-error@+2{{'device_global' attribute only applies to classes}} @@ -44,6 +121,12 @@ device_global same_name; // OK union [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] a_union; int main() { + // expected-note@+1{{in instantiation of function template specialization 'templFoo' requested here}} + templFoo(); + + // expected-note@+1{{in instantiation of function template specialization 'fooBar' requested here}} + fooBar(); + sycl::kernel_single_task([=]() { (void)glob; (void)static_glob; @@ -53,11 +136,7 @@ int main() { }); sycl::kernel_single_task([]() { - // expected-error@+1{{'device_global' variables must be static or declared at namespace scope}} + // expected-error@+1{{'device_global' variable must be a static data member or declared in global or namespace scope}} device_global non_static; - - // expect no error on non_const_static declaration if decorated with - // [[__sycl_detail__::global_variable_allowed]] - static device_global non_const_static; }); }