Skip to content

Commit 1265721

Browse files
schittirschittir
authored
[SYCL] Add the remaining diagnostics to device_global implementation (#5810)
This PR is a follow up to PR #5597 to implement the diagnostics not covered in #5597 As it stands currently, this PR includes implementation for cases described in this documentation and their templated versions - https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_device_global.asciidoc#restrictions-on-creating-device-global-objects except the shadow variable/namespace case. Specifically, it covers cases where device_global is declared as a private member inside a struct and within methods inside struct, their template versions/instantiation and device_global array.
1 parent c9f8938 commit 1265721

File tree

5 files changed

+173
-24
lines changed

5 files changed

+173
-24
lines changed

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

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -7219,9 +7219,11 @@ def warn_format_nonliteral : Warning<
72197219
InGroup<FormatNonLiteral>, DefaultIgnore;
72207220

72217221
def err_sycl_device_global_incorrect_scope : Error<
7222-
"'device_global' variables must be static or declared at namespace scope">;
7222+
"'device_global' variable must be a static data member or declared in global or namespace scope">;
72237223
def err_sycl_device_global_not_publicly_accessible: Error<
7224-
"'device_global' member variable %0 is not publicly accessible from namespace scope">;
7224+
"'device_global' member variable %0 should be publicly accessible from namespace scope">;
7225+
def err_sycl_device_global_array : Error<
7226+
"'device_global' array is not allowed">;
72257227

72267228
def err_unexpected_interface : Error<
72277229
"unexpected interface name %0: expected expression">;

‎clang/lib/Sema/SemaDecl.cpp

Lines changed: 25 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -7678,13 +7678,32 @@ NamedDecl *Sema::ActOnVariableDeclarator(
76787678
NewVD->setTSCSpec(TSCS);
76797679
}
76807680

7681-
// Global variables with types decorated with device_global attribute must be
7682-
// static if they are declared in SYCL device code.
76837681
if (getLangOpts().SYCLIsDevice) {
7684-
if (SCSpec != DeclSpec::SCS_static && !NewVD->hasGlobalStorage() &&
7685-
isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
7686-
NewVD->getType()))
7687-
Diag(D.getIdentifierLoc(), diag::err_sycl_device_global_incorrect_scope);
7682+
// device_global array is not allowed.
7683+
if (const ArrayType *AT = getASTContext().getAsArrayType(NewVD->getType()))
7684+
if (isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
7685+
AT->getElementType()))
7686+
Diag(NewVD->getLocation(), diag::err_sycl_device_global_array);
7687+
7688+
// Global variables with types decorated with device_global attribute must
7689+
// be static if they are declared in SYCL device code.
7690+
if (isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
7691+
NewVD->getType())) {
7692+
if (SCSpec == DeclSpec::SCS_static) {
7693+
const DeclContext *DC = NewVD->getDeclContext();
7694+
while (!DC->isTranslationUnit()) {
7695+
if (isa<FunctionDecl>(DC)) {
7696+
Diag(D.getIdentifierLoc(),
7697+
diag::err_sycl_device_global_incorrect_scope);
7698+
break;
7699+
}
7700+
DC = DC->getParent();
7701+
}
7702+
} else if (!NewVD->hasGlobalStorage()) {
7703+
Diag(D.getIdentifierLoc(),
7704+
diag::err_sycl_device_global_incorrect_scope);
7705+
}
7706+
}
76887707

76897708
// Static variables declared inside SYCL device code must be const or
76907709
// constexpr unless their types are decorated with global_variable_allowed

‎clang/lib/Sema/SemaDeclCXX.cpp

Lines changed: 18 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -3598,10 +3598,24 @@ Sema::ActOnCXXMemberDeclarator(Scope *S, AccessSpecifier AS, Declarator &D,
35983598
if (getLangOpts().SYCLIsDevice) {
35993599
if (auto Value = dyn_cast<ValueDecl>(Member)) {
36003600
if (isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
3601-
Value->getType()) &&
3602-
Value->getAccess() != AS_public) {
3603-
Diag(Loc, diag::err_sycl_device_global_not_publicly_accessible)
3604-
<< Value;
3601+
Value->getType())) {
3602+
if (Value->getAccess() == AS_private ||
3603+
Value->getAccess() == AS_protected) {
3604+
Diag(Loc, diag::err_sycl_device_global_not_publicly_accessible)
3605+
<< Value;
3606+
}
3607+
const DeclContext *DC = Member->getDeclContext();
3608+
while (!DC->isTranslationUnit()) {
3609+
if (auto Decl = dyn_cast<NamedDecl>(DC)) {
3610+
if (Decl->getAccess() == AS_private ||
3611+
Decl->getAccess() == AS_protected) {
3612+
Diag(Loc, diag::err_sycl_device_global_not_publicly_accessible)
3613+
<< Value;
3614+
break;
3615+
}
3616+
}
3617+
DC = DC->getParent();
3618+
}
36053619
}
36063620
}
36073621
}

‎clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

Lines changed: 36 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1614,14 +1614,38 @@ Decl *TemplateDeclInstantiator::VisitVarDecl(VarDecl *D,
16141614
// Only add this if we aren't instantiating a variable template. We'll end up
16151615
// adding the VarTemplateSpecializationDecl later.
16161616
if (!InstantiatingVarTemplate) {
1617-
SemaRef.addSyclVarDecl(Var);
1617+
if (SemaRef.getLangOpts().SYCLIsDevice &&
1618+
SemaRef.isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
1619+
Var->getType())) {
1620+
if (!Var->hasGlobalStorage())
1621+
SemaRef.Diag(D->getLocation(),
1622+
diag::err_sycl_device_global_incorrect_scope);
1623+
1624+
if (Var->getAccess() == AS_private || Var->getAccess() == AS_protected)
1625+
SemaRef.Diag(D->getLocation(),
1626+
diag::err_sycl_device_global_not_publicly_accessible)
1627+
<< Var;
1628+
1629+
if (Var->isStaticLocal()) {
1630+
const DeclContext *DC = Var->getDeclContext();
1631+
while (!DC->isTranslationUnit()) {
1632+
if (isa<FunctionDecl>(DC)) {
1633+
SemaRef.Diag(D->getLocation(),
1634+
diag::err_sycl_device_global_incorrect_scope);
1635+
break;
1636+
}
1637+
DC = DC->getParent();
1638+
}
1639+
}
1640+
}
16181641
if (const auto *SYCLDevice = Var->getAttr<SYCLDeviceAttr>()) {
16191642
if (!SemaRef.isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
16201643
Var->getType()))
16211644
SemaRef.Diag(SYCLDevice->getLoc(),
16221645
diag::err_sycl_attribute_not_device_global)
16231646
<< SYCLDevice;
16241647
}
1648+
SemaRef.addSyclVarDecl(Var);
16251649
}
16261650
return Var;
16271651
}
@@ -1711,6 +1735,17 @@ Decl *TemplateDeclInstantiator::VisitFieldDecl(FieldDecl *D) {
17111735

17121736
Field->setImplicit(D->isImplicit());
17131737
Field->setAccess(D->getAccess());
1738+
// Static members are not processed here, so error out if we have a device
1739+
// global without checking access modifier.
1740+
if (SemaRef.getLangOpts().SYCLIsDevice) {
1741+
if (SemaRef.isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
1742+
Field->getType())) {
1743+
SemaRef.Diag(D->getLocation(),
1744+
diag::err_sycl_device_global_incorrect_scope);
1745+
Field->setInvalidDecl();
1746+
return nullptr;
1747+
}
1748+
}
17141749
Owner->addDecl(Field);
17151750

17161751
return Field;

‎clang/test/SemaSYCL/device_global.cpp

Lines changed: 90 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -1,9 +1,10 @@
11
// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -verify %s
22
#include "Inputs/sycl.hpp"
33

4-
// Test cases below check for valid usage of device_global and
5-
// global_variable_allowed attributes, and that they are being correctly
6-
// generated in the AST.
4+
// Diagnostic tests for device_global and global_variable_allowed attribute.
5+
6+
// Test that there are no errors when variables of type device_global are
7+
// decorated with global_variable_allowed attribute appropriately.
78
using namespace sycl::ext::oneapi;
89

910
device_global<int> glob; // OK
@@ -18,19 +19,95 @@ device_global<char> Foo::d;
1819

1920
struct Baz {
2021
private:
21-
// expected-error@+1{{'device_global' member variable 'f' is not publicly accessible from namespace scope}}
22+
// expected-error@+1{{'device_global' member variable 'f' should be publicly accessible from namespace scope}}
2223
static device_global<int> f;
24+
25+
protected:
26+
// expected-error@+1{{'device_global' member variable 'g' should be publicly accessible from namespace scope}}
27+
static device_global<int> g;
2328
};
29+
2430
device_global<int> Baz::f;
2531

2632
device_global<int[4]> not_array; // OK
2733

34+
// expected-error@+1{{'device_global' array is not allowed}}
35+
device_global<int> array[4];
36+
2837
device_global<int> same_name; // OK
38+
2939
namespace foo {
3040
device_global<int> same_name; // OK
3141
}
32-
namespace {
33-
device_global<int> same_name; // OK
42+
43+
struct BBar {
44+
private:
45+
struct BarInsider {
46+
// expected-error@+1{{'device_global' member variable 'c' should be publicly accessible from namespace scope}}
47+
static device_global<float> c;
48+
};
49+
50+
protected:
51+
struct BarInsiderProtected {
52+
// expected-error@+1{{'device_global' member variable 'c' should be publicly accessible from namespace scope}}
53+
static device_global<float> c;
54+
};
55+
};
56+
57+
struct ABar {
58+
void method() {
59+
// expected-error@+1{{'device_global' variable must be a static data member or declared in global or namespace scope}}
60+
static device_global<float> c;
61+
}
62+
struct BarInsider {
63+
static device_global<float> c;
64+
void method() {
65+
// expected-error@+1{{'device_global' variable must be a static data member or declared in global or namespace scope}}
66+
static device_global<float> c;
67+
}
68+
};
69+
};
70+
71+
template <typename T> void fooBar() {
72+
// expected-error@+1{{'device_global' variable must be a static data member or declared in global or namespace scope}}
73+
static device_global<T> c;
74+
// expected-error@+1{{'device_global' variable must be a static data member or declared in global or namespace scope}}
75+
device_global<T> d;
76+
}
77+
78+
template <typename T> struct TS {
79+
private:
80+
// expected-error@+1 2{{'device_global' member variable 'a' should be publicly accessible from namespace scope}}
81+
static device_global<T> a;
82+
// expected-error@+1 2{{'device_global' variable must be a static data member or declared in global or namespace scope}}
83+
device_global<T> b;
84+
// expected-error@+2{{'device_global' member variable 'c' should be publicly accessible from namespace scope}}
85+
// expected-error@+1 2{{'device_global' variable must be a static data member or declared in global or namespace scope}}
86+
device_global<int> c;
87+
88+
public:
89+
static device_global<T> d;
90+
// expected-error@+1 2{{'device_global' variable must be a static data member or declared in global or namespace scope}}
91+
device_global<T> e;
92+
// expected-error@+1 2{{'device_global' variable must be a static data member or declared in global or namespace scope}}
93+
device_global<int> f;
94+
95+
protected:
96+
// expected-error@+1 2{{'device_global' member variable 'g' should be publicly accessible from namespace scope}}
97+
static device_global<T> g;
98+
// expected-error@+1 2{{'device_global' variable must be a static data member or declared in global or namespace scope}}
99+
device_global<T> h;
100+
// expected-error@+2{{'device_global' member variable 'i' should be publicly accessible from namespace scope}}
101+
// expected-error@+1 2{{'device_global' variable must be a static data member or declared in global or namespace scope}}
102+
device_global<int> i;
103+
};
104+
105+
// expected-note@+1{{in instantiation of template class 'TS<int>' requested here}}
106+
TS<int> AAAA;
107+
108+
//expected-note@+2{{in instantiation of template class 'TS<char>' requested here}}
109+
template <typename T> void templFoo () {
110+
TS<T> Var;
34111
}
35112

36113
// expected-error@+2{{'device_global' attribute only applies to classes}}
@@ -44,6 +121,12 @@ device_global<int> same_name; // OK
44121
union [[__sycl_detail__::device_global]] [[__sycl_detail__::global_variable_allowed]] a_union;
45122

46123
int main() {
124+
// expected-note@+1{{in instantiation of function template specialization 'templFoo<char>' requested here}}
125+
templFoo<char>();
126+
127+
// expected-note@+1{{in instantiation of function template specialization 'fooBar<int>' requested here}}
128+
fooBar<int>();
129+
47130
sycl::kernel_single_task<class KernelName1>([=]() {
48131
(void)glob;
49132
(void)static_glob;
@@ -53,11 +136,7 @@ int main() {
53136
});
54137

55138
sycl::kernel_single_task<class KernelName2>([]() {
56-
// expected-error@+1{{'device_global' variables must be static or declared at namespace scope}}
139+
// expected-error@+1{{'device_global' variable must be a static data member or declared in global or namespace scope}}
57140
device_global<int> non_static;
58-
59-
// expect no error on non_const_static declaration if decorated with
60-
// [[__sycl_detail__::global_variable_allowed]]
61-
static device_global<int> non_const_static;
62141
});
63142
}

0 commit comments

Comments
 (0)