diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index a05fe355d314d..416cdabf5de8d 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1197,6 +1197,19 @@ def SYCLDevice : InheritableAttr { let Documentation = [SYCLDeviceDocs]; } +def GlobalStorageNonLocalVar : SubsetSubjecthasGlobalStorage() && + !S->isLocalVarDeclOrParm()}], + "global variables">; + +def SYCLGlobalVar : InheritableAttr { + let Spellings = [GNU<"sycl_global_var">]; + let Subjects = SubjectList<[GlobalStorageNonLocalVar], ErrorDiag>; + let LangOpts = [SYCLIsDevice]; + // Only used internally by the SYCL implementation + let Documentation = [Undocumented]; +} + def SYCLKernel : InheritableAttr { let Spellings = [Clang<"sycl_kernel">]; let Subjects = SubjectList<[FunctionTmpl]>; diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 9d5c936c316d9..d24c88a690141 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -3402,6 +3402,8 @@ def warn_attribute_wrong_decl_type_str : Warning< "%0 attribute only applies to %1">, InGroup; def err_attribute_wrong_decl_type_str : Error< warn_attribute_wrong_decl_type_str.Text>; +def err_attribute_only_system_header : Error< + "%0 attribute only supported within a system header">; def warn_attribute_wrong_decl_type : Warning< "%0 attribute only applies to %select{" "functions" diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 548fa364d08d9..b8fc54f53da5d 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -5077,6 +5077,15 @@ static void handleSYCLDeviceIndirectlyCallableAttr(Sema &S, Decl *D, handleSimpleAttribute(S, D, AL); } +static void handleSYCLGlobalVarAttr(Sema &S, Decl *D, const ParsedAttr &AL) { + if (!S.Context.getSourceManager().isInSystemHeader(D->getLocation())) { + S.Diag(AL.getLoc(), diag::err_attribute_only_system_header) << AL; + return; + } + + handleSimpleAttribute(S, D, AL); +} + static void handleSYCLRegisterNumAttr(Sema &S, Decl *D, const ParsedAttr &AL) { if (!AL.checkExactlyNumArgs(S, 1)) return; @@ -9171,6 +9180,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, case ParsedAttr::AT_SYCLDeviceIndirectlyCallable: handleSYCLDeviceIndirectlyCallableAttr(S, D, AL); break; + case ParsedAttr::AT_SYCLGlobalVar: + handleSYCLGlobalVarAttr(S, D, AL); + break; case ParsedAttr::AT_SYCLRegisterNum: handleSYCLRegisterNumAttr(S, D, AL); break; diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index cf3bc4300247b..71f3872e51bee 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -220,12 +220,15 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, ExprEvalContexts.empty() || (!isUnevaluatedContext() && !isConstantEvaluated()); bool IsEsimdPrivateGlobal = isSYCLEsimdPrivateGlobal(VD); - if (IsRuntimeEvaluated && !IsConst && VD->getStorageClass() == SC_Static) + if (IsRuntimeEvaluated && !IsConst && + VD->getStorageClass() == SC_Static && + !VD->hasAttr()) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelNonConstStaticDataVariable; - // Non-const globals are allowed for SYCL explicit SIMD. + // Non-const globals are allowed for SYCL explicit SIMD or with the + // SYCLGlobalVar attribute. else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && - VD->hasGlobalStorage()) + VD->hasGlobalStorage() && !VD->hasAttr()) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelGlobalVariable; // ESIMD globals cannot be used in a SYCL context. diff --git a/clang/test/SemaSYCL/attr-syclglobalvar.cpp b/clang/test/SemaSYCL/attr-syclglobalvar.cpp new file mode 100644 index 0000000000000..968e556d5a72d --- /dev/null +++ b/clang/test/SemaSYCL/attr-syclglobalvar.cpp @@ -0,0 +1,150 @@ +// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only %s + +#include "Inputs/sycl.hpp" + +# 5 "header.hpp" 1 3 // Simulate a system #include to enter new file named header.hpp at line 5 + +#define SYCLGLOBALVAR_ATTR_MACRO __attribute__((sycl_global_var)) + +__attribute__((sycl_global_var)) int HppGlobalWithAttribute; + +__attribute__((sycl_global_var)) extern int HppExternGlobalWithAttribute; + +namespace NS { + __attribute__((sycl_global_var)) int HppNSGlobalWithAttribute; +} + +struct HppS { + __attribute__((sycl_global_var)) static int StaticMember; + + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int InstanceMember; +}; +int HppS::StaticMember = 0; + +__attribute__((sycl_global_var)) HppS HppGlobalStruct; + +__attribute__((sycl_global_var)) static HppS HppStaticGlobal; + +static union { + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int HppAnonymousStaticUnionInstanceMember; +}; + +// expected-error@+1 {{attribute takes no arguments}} +__attribute__((sycl_global_var(42))) int HppGlobalWithAttributeArg; + +template struct HppStructTemplate { + __attribute__((sycl_global_var)) static T StaticMember; + + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int InstanceMember; +}; + +SYCLGLOBALVAR_ATTR_MACRO int HppGlobalWithAttrMacro; + +int HppGlobalNoAttribute; + +// expected-error@+1 {{attribute only applies to global variables}} +__attribute__((sycl_global_var)) void HppF( + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int Param +) { + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) static int StaticLocalVar; + + // expected-error@+1 {{attribute only applies to global variables}} + __attribute__((sycl_global_var)) int Local; + + cl::sycl::kernel_single_task([=] () { + (void)HppGlobalWithAttribute; // ok + (void)HppExternGlobalWithAttribute; // ok + (void)NS::HppNSGlobalWithAttribute; // ok + (void)HppS::StaticMember; // ok + (void)HppGlobalStruct.InstanceMember; // ok + (void)HppStaticGlobal.InstanceMember; // ok + (void)HppAnonymousStaticUnionInstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}} + (void)HppGlobalWithAttributeArg; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)HppStructTemplate::StaticMember; // ok + (void)HppGlobalWithAttrMacro; // ok + (void)HppGlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@Inputs/sycl.hpp:* {{called by}} + }); +} + +# 74 "header.hpp" 2 // Return from the simulated #include (with the last line number of the "header.hpp" file) + +// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}} +__attribute__((sycl_global_var)) int CppGlobalWithAttribute; + +// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}} +__attribute__((sycl_global_var)) extern int CppExternGlobalWithAttribute; + +namespace NS { + // expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}} + __attribute__((sycl_global_var)) int CppNSGlobalWithAttribute; +} + +struct CppS { + // expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}} + __attribute__((sycl_global_var)) static int StaticMember; + + // expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}} + __attribute__((sycl_global_var)) int InstanceMember; +}; +int CppS::StaticMember = 0; + +// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}} +__attribute__((sycl_global_var)) CppS CppGlobalStruct; + +// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}} +__attribute__((sycl_global_var)) static CppS CppStaticGlobal; + +static union { + // expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}} + __attribute__((sycl_global_var)) int CppAnonymousStaticUnionInstanceMember; +}; + +// expected-error@+1 {{attribute takes no arguments}} +__attribute__((sycl_global_var(42))) int CppGlobalWithAttributeArg; + +// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}} +__attribute__((sycl_global_var)) HppStructTemplate CppGlobalTemplateStructWithAttribute; +HppStructTemplate CppGlobalTemplateStructNoAttribute; + +// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}} +SYCLGLOBALVAR_ATTR_MACRO int CppGlobalWithAttrMacro; + +int GlobalNoAttribute; + +// expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}} +__attribute__((sycl_global_var)) void F( + // expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}} + __attribute__((sycl_global_var)) int Param +) { + // expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}} + __attribute__((sycl_global_var)) static int StaticLocalVar; + + // expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}} + __attribute__((sycl_global_var)) int Local; + + cl::sycl::kernel_single_task([=] () { + (void)HppGlobalWithAttribute; // ok + (void)CppGlobalWithAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)HppExternGlobalWithAttribute; // ok + (void)CppExternGlobalWithAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)NS::HppNSGlobalWithAttribute; // ok + (void)NS::CppNSGlobalWithAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)HppS::StaticMember; // ok + (void)CppS::StaticMember; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)HppGlobalStruct.InstanceMember; // ok + (void)CppGlobalStruct.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)HppStaticGlobal.InstanceMember; // ok + (void)CppStaticGlobal.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}} + (void)CppAnonymousStaticUnionInstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}} + (void)CppGlobalWithAttributeArg; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)HppStructTemplate::StaticMember; // ok + (void)CppGlobalTemplateStructWithAttribute.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)CppGlobalTemplateStructNoAttribute.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const global variable}} + (void)GlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@Inputs/sycl.hpp:* {{called by}} + }); +}