Skip to content

Commit 31bdbbf

Browse files
jtmott-inteljtmott-intel
authored
[SYCL][Clang] Add sycl_global_var attribute (#3746)
Normally global variables are disallowed within kernels, but the presence of this new sycl_global_var attribute will cause Sema to allow that particular global variable.
1 parent 5759e2a commit 31bdbbf

File tree

5 files changed

+183
-3
lines changed

5 files changed

+183
-3
lines changed

‎clang/include/clang/Basic/Attr.td

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1206,6 +1206,19 @@ def SYCLDevice : InheritableAttr {
12061206
let Documentation = [SYCLDeviceDocs];
12071207
}
12081208

1209+
def GlobalStorageNonLocalVar : SubsetSubject<Var,
1210+
[{S->hasGlobalStorage() &&
1211+
!S->isLocalVarDeclOrParm()}],
1212+
"global variables">;
1213+
1214+
def SYCLGlobalVar : InheritableAttr {
1215+
let Spellings = [GNU<"sycl_global_var">];
1216+
let Subjects = SubjectList<[GlobalStorageNonLocalVar], ErrorDiag>;
1217+
let LangOpts = [SYCLIsDevice];
1218+
// Only used internally by the SYCL implementation
1219+
let Documentation = [Undocumented];
1220+
}
1221+
12091222
def SYCLKernel : InheritableAttr {
12101223
let Spellings = [Clang<"sycl_kernel">];
12111224
let Subjects = SubjectList<[FunctionTmpl]>;

‎clang/include/clang/Basic/DiagnosticSemaKinds.td

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3424,6 +3424,8 @@ def warn_attribute_wrong_decl_type_str : Warning<
34243424
"%0 attribute only applies to %1">, InGroup<IgnoredAttributes>;
34253425
def err_attribute_wrong_decl_type_str : Error<
34263426
warn_attribute_wrong_decl_type_str.Text>;
3427+
def err_attribute_only_system_header : Error<
3428+
"%0 attribute only supported within a system header">;
34273429
def warn_attribute_wrong_decl_type : Warning<
34283430
"%0 attribute only applies to %select{"
34293431
"functions"

‎clang/lib/Sema/SemaDeclAttr.cpp

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -5362,6 +5362,15 @@ static void handleSYCLDeviceIndirectlyCallableAttr(Sema &S, Decl *D,
53625362
handleSimpleAttribute<SYCLDeviceIndirectlyCallableAttr>(S, D, AL);
53635363
}
53645364

5365+
static void handleSYCLGlobalVarAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
5366+
if (!S.Context.getSourceManager().isInSystemHeader(D->getLocation())) {
5367+
S.Diag(AL.getLoc(), diag::err_attribute_only_system_header) << AL;
5368+
return;
5369+
}
5370+
5371+
handleSimpleAttribute<SYCLGlobalVarAttr>(S, D, AL);
5372+
}
5373+
53655374
static void handleSYCLRegisterNumAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
53665375
if (!AL.checkExactlyNumArgs(S, 1))
53675376
return;
@@ -9503,6 +9512,9 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
95039512
case ParsedAttr::AT_SYCLDeviceIndirectlyCallable:
95049513
handleSYCLDeviceIndirectlyCallableAttr(S, D, AL);
95059514
break;
9515+
case ParsedAttr::AT_SYCLGlobalVar:
9516+
handleSYCLGlobalVarAttr(S, D, AL);
9517+
break;
95069518
case ParsedAttr::AT_SYCLRegisterNum:
95079519
handleSYCLRegisterNumAttr(S, D, AL);
95089520
break;

‎clang/lib/Sema/SemaExpr.cpp

Lines changed: 6 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -223,12 +223,15 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
223223
ExprEvalContexts.empty() ||
224224
(!isUnevaluatedContext() && !isConstantEvaluated());
225225
bool IsEsimdPrivateGlobal = isSYCLEsimdPrivateGlobal(VD);
226-
if (IsRuntimeEvaluated && !IsConst && VD->getStorageClass() == SC_Static)
226+
if (IsRuntimeEvaluated && !IsConst &&
227+
VD->getStorageClass() == SC_Static &&
228+
!VD->hasAttr<SYCLGlobalVarAttr>())
227229
SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict)
228230
<< Sema::KernelNonConstStaticDataVariable;
229-
// Non-const globals are allowed for SYCL explicit SIMD.
231+
// Non-const globals are allowed for SYCL explicit SIMD or with the
232+
// SYCLGlobalVar attribute.
230233
else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst &&
231-
VD->hasGlobalStorage())
234+
VD->hasGlobalStorage() && !VD->hasAttr<SYCLGlobalVarAttr>())
232235
SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict)
233236
<< Sema::KernelGlobalVariable;
234237
// ESIMD globals cannot be used in a SYCL context.

‎clang/test/SemaSYCL/attr-syclglobalvar.cpp

Lines changed: 150 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,150 @@
1+
// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only %s
2+
3+
#include "Inputs/sycl.hpp"
4+
5+
# 5 "header.hpp" 1 3 // Simulate a system #include to enter new file named header.hpp at line 5
6+
7+
#define SYCLGLOBALVAR_ATTR_MACRO __attribute__((sycl_global_var))
8+
9+
__attribute__((sycl_global_var)) int HppGlobalWithAttribute;
10+
11+
__attribute__((sycl_global_var)) extern int HppExternGlobalWithAttribute;
12+
13+
namespace NS {
14+
__attribute__((sycl_global_var)) int HppNSGlobalWithAttribute;
15+
}
16+
17+
struct HppS {
18+
__attribute__((sycl_global_var)) static int StaticMember;
19+
20+
// expected-error@+1 {{attribute only applies to global variables}}
21+
__attribute__((sycl_global_var)) int InstanceMember;
22+
};
23+
int HppS::StaticMember = 0;
24+
25+
__attribute__((sycl_global_var)) HppS HppGlobalStruct;
26+
27+
__attribute__((sycl_global_var)) static HppS HppStaticGlobal;
28+
29+
static union {
30+
// expected-error@+1 {{attribute only applies to global variables}}
31+
__attribute__((sycl_global_var)) int HppAnonymousStaticUnionInstanceMember;
32+
};
33+
34+
// expected-error@+1 {{attribute takes no arguments}}
35+
__attribute__((sycl_global_var(42))) int HppGlobalWithAttributeArg;
36+
37+
template<typename T> struct HppStructTemplate {
38+
__attribute__((sycl_global_var)) static T StaticMember;
39+
40+
// expected-error@+1 {{attribute only applies to global variables}}
41+
__attribute__((sycl_global_var)) int InstanceMember;
42+
};
43+
44+
SYCLGLOBALVAR_ATTR_MACRO int HppGlobalWithAttrMacro;
45+
46+
int HppGlobalNoAttribute;
47+
48+
// expected-error@+1 {{attribute only applies to global variables}}
49+
__attribute__((sycl_global_var)) void HppF(
50+
// expected-error@+1 {{attribute only applies to global variables}}
51+
__attribute__((sycl_global_var)) int Param
52+
) {
53+
// expected-error@+1 {{attribute only applies to global variables}}
54+
__attribute__((sycl_global_var)) static int StaticLocalVar;
55+
56+
// expected-error@+1 {{attribute only applies to global variables}}
57+
__attribute__((sycl_global_var)) int Local;
58+
59+
cl::sycl::kernel_single_task<class kernel_name>([=] () {
60+
(void)HppGlobalWithAttribute; // ok
61+
(void)HppExternGlobalWithAttribute; // ok
62+
(void)NS::HppNSGlobalWithAttribute; // ok
63+
(void)HppS::StaticMember; // ok
64+
(void)HppGlobalStruct.InstanceMember; // ok
65+
(void)HppStaticGlobal.InstanceMember; // ok
66+
(void)HppAnonymousStaticUnionInstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}}
67+
(void)HppGlobalWithAttributeArg; // expected-error {{SYCL kernel cannot use a non-const global variable}}
68+
(void)HppStructTemplate<int>::StaticMember; // ok
69+
(void)HppGlobalWithAttrMacro; // ok
70+
(void)HppGlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@Inputs/sycl.hpp:* {{called by}}
71+
});
72+
}
73+
74+
# 74 "header.hpp" 2 // Return from the simulated #include (with the last line number of the "header.hpp" file)
75+
76+
// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}}
77+
__attribute__((sycl_global_var)) int CppGlobalWithAttribute;
78+
79+
// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}}
80+
__attribute__((sycl_global_var)) extern int CppExternGlobalWithAttribute;
81+
82+
namespace NS {
83+
// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}}
84+
__attribute__((sycl_global_var)) int CppNSGlobalWithAttribute;
85+
}
86+
87+
struct CppS {
88+
// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}}
89+
__attribute__((sycl_global_var)) static int StaticMember;
90+
91+
// expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}}
92+
__attribute__((sycl_global_var)) int InstanceMember;
93+
};
94+
int CppS::StaticMember = 0;
95+
96+
// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}}
97+
__attribute__((sycl_global_var)) CppS CppGlobalStruct;
98+
99+
// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}}
100+
__attribute__((sycl_global_var)) static CppS CppStaticGlobal;
101+
102+
static union {
103+
// expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}}
104+
__attribute__((sycl_global_var)) int CppAnonymousStaticUnionInstanceMember;
105+
};
106+
107+
// expected-error@+1 {{attribute takes no arguments}}
108+
__attribute__((sycl_global_var(42))) int CppGlobalWithAttributeArg;
109+
110+
// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}}
111+
__attribute__((sycl_global_var)) HppStructTemplate<int> CppGlobalTemplateStructWithAttribute;
112+
HppStructTemplate<int> CppGlobalTemplateStructNoAttribute;
113+
114+
// expected-error@+1 {{'sycl_global_var' attribute only supported within a system header}}
115+
SYCLGLOBALVAR_ATTR_MACRO int CppGlobalWithAttrMacro;
116+
117+
int GlobalNoAttribute;
118+
119+
// expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}}
120+
__attribute__((sycl_global_var)) void F(
121+
// expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}}
122+
__attribute__((sycl_global_var)) int Param
123+
) {
124+
// expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}}
125+
__attribute__((sycl_global_var)) static int StaticLocalVar;
126+
127+
// expected-error@+1 {{'sycl_global_var' attribute only applies to global variables}}
128+
__attribute__((sycl_global_var)) int Local;
129+
130+
cl::sycl::kernel_single_task<class kernel_name>([=] () {
131+
(void)HppGlobalWithAttribute; // ok
132+
(void)CppGlobalWithAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}}
133+
(void)HppExternGlobalWithAttribute; // ok
134+
(void)CppExternGlobalWithAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}}
135+
(void)NS::HppNSGlobalWithAttribute; // ok
136+
(void)NS::CppNSGlobalWithAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}}
137+
(void)HppS::StaticMember; // ok
138+
(void)CppS::StaticMember; // expected-error {{SYCL kernel cannot use a non-const global variable}}
139+
(void)HppGlobalStruct.InstanceMember; // ok
140+
(void)CppGlobalStruct.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const global variable}}
141+
(void)HppStaticGlobal.InstanceMember; // ok
142+
(void)CppStaticGlobal.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}}
143+
(void)CppAnonymousStaticUnionInstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}}
144+
(void)CppGlobalWithAttributeArg; // expected-error {{SYCL kernel cannot use a non-const global variable}}
145+
(void)HppStructTemplate<int>::StaticMember; // ok
146+
(void)CppGlobalTemplateStructWithAttribute.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const global variable}}
147+
(void)CppGlobalTemplateStructNoAttribute.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const global variable}}
148+
(void)GlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@Inputs/sycl.hpp:* {{called by}}
149+
});
150+
}

0 commit comments

Comments
 (0)