Skip to content

Commit 8b958f6

Browse files
FznamznonFznamznon
authored
[SYCL] Add support for SYCL_EXTERNAL on device_global variables (#5963)
Device global spec declares that it is possible to apply `SYCL_EXTERNAL` macro to `device_global` variables. This patch modifies `sycl_device` attribute which is used to implement `SYCL_EXTERNAL` macro so it is possible to apply `SYCL_EXTERNAL` macro to `device_global` variables and it is not allowed to apply it to any other variables.
1 parent 770f540 commit 8b958f6

File tree

9 files changed

+122
-15
lines changed

9 files changed

+122
-15
lines changed

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

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1247,18 +1247,18 @@ def CUDAShared : InheritableAttr {
12471247
}
12481248
def : MutualExclusions<[CUDAConstant, CUDAShared, HIPManaged]>;
12491249

1250+
def GlobalStorageNonLocalVar : SubsetSubject<Var,
1251+
[{S->hasGlobalStorage() &&
1252+
!S->isLocalVarDeclOrParm()}],
1253+
"global variables">;
1254+
12501255
def SYCLDevice : InheritableAttr {
12511256
let Spellings = [GNU<"sycl_device">];
1252-
let Subjects = SubjectList<[Function]>;
1257+
let Subjects = SubjectList<[Function, GlobalStorageNonLocalVar]>;
12531258
let LangOpts = [SYCLIsDevice, SilentlyIgnoreSYCLIsHost];
12541259
let Documentation = [SYCLDeviceDocs];
12551260
}
12561261

1257-
def GlobalStorageNonLocalVar : SubsetSubject<Var,
1258-
[{S->hasGlobalStorage() &&
1259-
!S->isLocalVarDeclOrParm()}],
1260-
"global variables">;
1261-
12621262
def SYCLGlobalVar : InheritableAttr {
12631263
let Spellings = [GNU<"sycl_global_var">];
12641264
let Subjects = SubjectList<[GlobalStorageNonLocalVar], ErrorDiag>;

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

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -11727,6 +11727,9 @@ def err_sycl_restrict : Error<
1172711727
"|use a const static or global variable that is neither zero-initialized "
1172811728
"nor constant-initialized"
1172911729
"}0">;
11730+
def err_sycl_external_global : Error<
11731+
"invalid reference to 'device_global' variable; external 'device_global'"
11732+
" variable must be marked with SYCL_EXTERNAL macro">;
1173011733
def warn_sycl_kernel_too_big_args : Warning<
1173111734
"size of kernel arguments (%0 bytes) may exceed the supported maximum "
1173211735
"of %1 bytes on some devices">, InGroup<SyclStrict>, ShowInSystemHeader;
@@ -11758,9 +11761,12 @@ def err_sycl_function_attribute_mismatch : Error<
1175811761
"SYCL kernel without %0 attribute can't call a function with this attribute">;
1175911762
def err_sycl_x_y_z_arguments_must_be_one : Error<
1176011763
"all %0 attribute arguments must be '1' when the %1 attribute argument is '0'">;
11761-
def err_sycl_attribute_internal_function
11764+
def err_sycl_attribute_internal_decl
1176211765
: Error<"%0 attribute cannot be applied to a "
11763-
"static function or function in an anonymous namespace">;
11766+
"static %select{function|variable}1 or %select{function|variable}1 "
11767+
"in an anonymous namespace">;
11768+
def err_sycl_attribute_not_device_global
11769+
: Error<"%0 attribute can only be applied to 'device_global' variables">;
1176411770
def err_sycl_compiletime_property_duplication : Error<
1176511771
"can't apply %0 property twice to the same accessor">;
1176611772
def err_sycl_invalid_property_list_param_number : Error<

‎clang/lib/Sema/Sema.cpp

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1715,6 +1715,11 @@ class DeferredDiagnosticsEmitter
17151715
<< Sema::KernelConstStaticVariable;
17161716
return;
17171717
}
1718+
if (!VD->hasInit() &&
1719+
S.isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
1720+
VD->getType()) &&
1721+
!VD->hasAttr<SYCLDeviceAttr>())
1722+
S.Diag(Loc, diag::err_sycl_external_global);
17181723
}
17191724
}
17201725
if (isa<VarDecl>(D))

‎clang/lib/Sema/SemaDeclAttr.cpp

Lines changed: 18 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -6014,20 +6014,34 @@ static void handleOptimizeNoneAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
60146014
}
60156015

60166016
static void handleSYCLDeviceAttr(Sema &S, Decl *D, const ParsedAttr &AL) {
6017-
auto *FD = cast<FunctionDecl>(D);
6018-
if (!FD->isExternallyVisible()) {
6019-
S.Diag(AL.getLoc(), diag::err_sycl_attribute_internal_function) << AL;
6017+
auto *ND = cast<NamedDecl>(D);
6018+
if (!ND->isExternallyVisible()) {
6019+
S.Diag(AL.getLoc(), diag::err_sycl_attribute_internal_decl)
6020+
<< AL << !isa<FunctionDecl>(ND);
60206021
return;
60216022
}
60226023

6024+
if (auto *VD = dyn_cast<VarDecl>(D)) {
6025+
QualType VarType = VD->getType();
6026+
// Diagnose only for non-dependent types since dependent type don't have
6027+
// attributes applied on them ATM.
6028+
if (!VarType->isDependentType() &&
6029+
!S.isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
6030+
VD->getType())) {
6031+
S.Diag(AL.getLoc(), diag::err_sycl_attribute_not_device_global) << AL;
6032+
return;
6033+
}
6034+
}
6035+
60236036
handleSimpleAttribute<SYCLDeviceAttr>(S, D, AL);
60246037
}
60256038

60266039
static void handleSYCLDeviceIndirectlyCallableAttr(Sema &S, Decl *D,
60276040
const ParsedAttr &AL) {
60286041
auto *FD = cast<FunctionDecl>(D);
60296042
if (!FD->isExternallyVisible()) {
6030-
S.Diag(AL.getLoc(), diag::err_sycl_attribute_internal_function) << AL;
6043+
S.Diag(AL.getLoc(), diag::err_sycl_attribute_internal_decl)
6044+
<< AL << /*function*/ 0;
60316045
return;
60326046
}
60336047

‎clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

Lines changed: 9 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1613,8 +1613,16 @@ Decl *TemplateDeclInstantiator::VisitVarDecl(VarDecl *D,
16131613

16141614
// Only add this if we aren't instantiating a variable template. We'll end up
16151615
// adding the VarTemplateSpecializationDecl later.
1616-
if (!InstantiatingVarTemplate)
1616+
if (!InstantiatingVarTemplate) {
16171617
SemaRef.addSyclVarDecl(Var);
1618+
if (const auto *SYCLDevice = Var->getAttr<SYCLDeviceAttr>()) {
1619+
if (!SemaRef.isTypeDecoratedWithDeclAttribute<SYCLDeviceGlobalAttr>(
1620+
Var->getType()))
1621+
SemaRef.Diag(SYCLDevice->getLoc(),
1622+
diag::err_sycl_attribute_not_device_global)
1623+
<< SYCLDevice;
1624+
}
1625+
}
16181626
return Var;
16191627
}
16201628

‎clang/test/CodeGenSYCL/device_global.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,12 +10,14 @@ using namespace sycl;
1010
queue q;
1111

1212
device_global<int> A;
13+
SYCL_EXTERNAL device_global<int> AExt;
1314
static device_global<int> B;
1415

1516
struct Foo {
1617
static device_global<int> C;
1718
};
1819
device_global<int> Foo::C;
20+
// CHECK: @AExt = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[AEXT_ATTRS:[0-9]+]]
1921
// CHECK: @A = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[A_ATTRS:[0-9]+]]
2022
// CHECK: @_ZL1B = internal addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[B_ATTRS:[0-9]+]]
2123
// CHECK: @_ZN3Foo1CE = addrspace(1) global %"class.sycl::_V1::ext::oneapi::device_global" zeroinitializer, align 8 #[[C_ATTRS:[0-9]+]]
@@ -95,6 +97,7 @@ void bar() {
9597
// CHECK-SAME: @_ZL1B
9698
// CHECK-SAME: @_ZN12_GLOBAL__N_19same_nameE
9799

100+
// CHECK: attributes #[[AEXT_ATTRS]] = { "sycl-unique-id"="_Z4AExt" }
98101
// CHECK: attributes #[[A_ATTRS]] = { "sycl-unique-id"="_Z1A" }
99102
// CHECK: attributes #[[B_ATTRS]] = { "sycl-unique-id"="THE_PREFIX____ZL1B" }
100103
// CHECK: attributes #[[C_ATTRS]] = { "sycl-unique-id"="_ZN3Foo1CE" }

‎clang/test/Misc/pragma-attribute-supported-attributes-list.test

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -164,7 +164,6 @@
164164
// CHECK-NEXT: SYCLAddIRAttributesFunction (SubjectMatchRule_function)
165165
// CHECK-NEXT: SYCLAddIRAttributesGlobalVariable (SubjectMatchRule_record)
166166
// CHECK-NEXT: SYCLAddIRAttributesKernelParameter (SubjectMatchRule_variable_is_parameter)
167-
// CHECK-NEXT: SYCLDevice (SubjectMatchRule_function)
168167
// CHECK-NEXT: SYCLDeviceGlobal (SubjectMatchRule_record)
169168
// CHECK-NEXT: SYCLDeviceHas (SubjectMatchRule_function)
170169
// CHECK-NEXT: SYCLDeviceIndirectlyCallable (SubjectMatchRule_function)

‎clang/test/SemaSYCL/device_global_external.cpp

Lines changed: 72 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,72 @@
1+
// RUN: %clang_cc1 -fsycl-is-device -std=c++17 -sycl-std=2020 -verify %s
2+
#include "Inputs/sycl.hpp"
3+
4+
// Tests that SYCL_EXTERNAL can be applied to device_global variables, and cannot be applied to other variables.
5+
using namespace sycl::ext::oneapi;
6+
7+
SYCL_EXTERNAL device_global<int> glob;
8+
// expected-error@+1{{'sycl_device' attribute cannot be applied to a static variable or variable in an anonymous namespace}}
9+
SYCL_EXTERNAL static device_global<float> static_glob;
10+
11+
namespace foo {
12+
SYCL_EXTERNAL device_global<int> same_name;
13+
}
14+
15+
struct RandomStruct {
16+
int M;
17+
};
18+
19+
// expected-error@+1{{'sycl_device' attribute can only be applied to 'device_global' variables}}
20+
SYCL_EXTERNAL RandomStruct S;
21+
22+
namespace {
23+
// expected-error@+1{{'sycl_device' attribute cannot be applied to a static variable or variable in an anonymous namespace}}
24+
SYCL_EXTERNAL device_global<int> same_name;
25+
} // namespace
26+
27+
// expected-error@+1{{'sycl_device' attribute can only be applied to 'device_global' variables}}
28+
SYCL_EXTERNAL int AAA;
29+
30+
struct B {
31+
SYCL_EXTERNAL static device_global<int> Member;
32+
};
33+
34+
void foofoo() {
35+
// expected-warning@+1{{'sycl_device' attribute only applies to functions and global variables}}
36+
SYCL_EXTERNAL RandomStruct S;
37+
// expected-warning@+1{{'sycl_device' attribute only applies to functions and global variables}}
38+
SYCL_EXTERNAL int A;
39+
}
40+
41+
template <typename T> struct NonDevGlob {
42+
};
43+
44+
template <typename T> struct TS {
45+
SYCL_EXTERNAL static device_global<T> D;
46+
// expected-error@+1{{'sycl_device' attribute can only be applied to 'device_global' variables}}
47+
SYCL_EXTERNAL static NonDevGlob<T> ND;
48+
};
49+
50+
// expected-note@+1 {{in instantiation of template class 'TS<int>' requested here}}
51+
TS<int> A;
52+
53+
struct [[__sycl_detail__::global_variable_allowed]] GlobAllowedOnly {
54+
};
55+
56+
// expected-error@+1{{'sycl_device' attribute can only be applied to 'device_global' variables}}
57+
SYCL_EXTERNAL GlobAllowedOnly GAO;
58+
59+
60+
SYCL_EXTERNAL extern device_global<int> Good;
61+
extern device_global<int> Bad;
62+
63+
int main() {
64+
sycl::kernel_single_task<class KernelName1>([=]() {
65+
Good.get();
66+
// expected-error@+1 {{invalid reference to 'device_global' variable; external 'device_global' variable must be marked with SYCL_EXTERNAL macro}}
67+
Bad.get();
68+
69+
(void)GAO;
70+
});
71+
return 0;
72+
}

‎clang/test/SemaSYCL/sycl-device.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -6,7 +6,7 @@
66

77
#ifdef SYCL
88

9-
__attribute__((sycl_device)) // expected-warning {{'sycl_device' attribute only applies to functions}}
9+
__attribute__((sycl_device)) // expected-error {{'sycl_device' attribute can only be applied to 'device_global' variables}}
1010
int N;
1111

1212
__attribute__((sycl_device(3))) // expected-error {{'sycl_device' attribute takes no arguments}}

0 commit comments

Comments
 (0)