Skip to content

Commit 7a2c8df

Browse files
baderbader
committed
[SYCL] Re-use OpenCL address space attributes for SYCL
Today we re-use OpenCL parsed attributes, but have separate SYCL address space semantic attributes as current implementation of OpenCL semantics breaks valid C++. This patch enables re-use of OpenCL semantic attributes by allowing conversions between types qualified with OpenCL address spaces and type w/o address space qualifiers. Clang compiler (almost) always adds address space qualifiers in OpenCL mode, so it should not affect OpenCL mode. NOTE: this change also disables implicit conversion between the unqualified types and types qualified with `__attribute__((address_space(N)))`, enabled by one of the previous SYCL patches. Signed-off-by: Alexey Bader <[email protected]>
1 parent 54dddb4 commit 7a2c8df

File tree

17 files changed

+50
-200
lines changed

17 files changed

+50
-200
lines changed

‎clang/include/clang/AST/Type.h

Lines changed: 6 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -486,15 +486,12 @@ class Qualifiers {
486486
/// Returns true if the address space in these qualifiers is equal to or
487487
/// a superset of the address space in the argument qualifiers.
488488
bool isAddressSpaceSupersetOf(Qualifiers other) const {
489-
490-
return
491-
isAddressSpaceSupersetOf(getAddressSpace(), other.getAddressSpace()) ||
492-
(!hasAddressSpace() &&
493-
(other.getAddressSpace() == LangAS::sycl_private ||
494-
other.getAddressSpace() == LangAS::sycl_local ||
495-
other.getAddressSpace() == LangAS::sycl_global ||
496-
other.getAddressSpace() == LangAS::sycl_constant ||
497-
other.getAddressSpace() == LangAS::sycl_generic));
489+
return isAddressSpaceSupersetOf(getAddressSpace(),
490+
other.getAddressSpace()) ||
491+
(!hasAddressSpace() &&
492+
(other.getAddressSpace() == LangAS::opencl_private ||
493+
other.getAddressSpace() == LangAS::opencl_local ||
494+
other.getAddressSpace() == LangAS::opencl_global));
498495
}
499496

500497
/// Determines if these qualifiers compatibly include another set.

‎clang/include/clang/Basic/AddressSpaces.h

Lines changed: 0 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -42,14 +42,6 @@ enum class LangAS : unsigned {
4242
cuda_constant,
4343
cuda_shared,
4444

45-
sycl_global,
46-
sycl_local,
47-
sycl_constant,
48-
sycl_private,
49-
// Likely never used, but useful in the future to reserve the spot in the
50-
// enum.
51-
sycl_generic,
52-
5345
// Pointer size and extension address spaces.
5446
ptr32_sptr,
5547
ptr32_uptr,

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

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -10292,8 +10292,6 @@ def err_builtin_launder_invalid_arg : Error<
1029210292
"'__builtin_launder' is not allowed">;
1029310293

1029410294
// SYCL-specific diagnostics
10295-
def err_sycl_attribute_address_space_invalid : Error<
10296-
"address space is outside the valid range of values">;
1029710295
def err_sycl_kernel_name_class_not_top_level : Error<
1029810296
"kernel name class and its template argument classes' declarations can only "
1029910297
"nest in a namespace: %0">;

‎clang/include/clang/Sema/ParsedAttr.h

Lines changed: 0 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -534,24 +534,6 @@ class ParsedAttr final
534534
}
535535
}
536536

537-
/// If this is an OpenCL addr space attribute returns its SYCL representation
538-
/// in LangAS, otherwise returns default addr space.
539-
LangAS asSYCLLangAS() const {
540-
switch (getKind()) {
541-
case ParsedAttr::AT_OpenCLConstantAddressSpace:
542-
return LangAS::sycl_constant;
543-
case ParsedAttr::AT_OpenCLGlobalAddressSpace:
544-
return LangAS::sycl_global;
545-
case ParsedAttr::AT_OpenCLLocalAddressSpace:
546-
return LangAS::sycl_local;
547-
case ParsedAttr::AT_OpenCLPrivateAddressSpace:
548-
return LangAS::sycl_private;
549-
case ParsedAttr::AT_OpenCLGenericAddressSpace:
550-
default:
551-
return LangAS::Default;
552-
}
553-
}
554-
555537
AttributeCommonInfo::Kind getKind() const { return getParsedKind(); }
556538
};
557539

‎clang/lib/AST/ASTContext.cpp

Lines changed: 0 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -836,11 +836,6 @@ static const LangASMap *getAddressSpaceMap(const TargetInfo &T,
836836
5, // cuda_device
837837
6, // cuda_constant
838838
7, // cuda_shared
839-
1, // sycl_global
840-
3, // sycl_local
841-
2, // sycl_constant
842-
0, // sycl_private
843-
4, // sycl_generic
844839
8, // ptr32_sptr
845840
9, // ptr32_uptr
846841
10 // ptr64

‎clang/lib/AST/TypePrinter.cpp

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1792,16 +1792,12 @@ std::string Qualifiers::getAddrSpaceAsString(LangAS AS) {
17921792
case LangAS::Default:
17931793
return "";
17941794
case LangAS::opencl_global:
1795-
case LangAS::sycl_global:
17961795
return "__global";
17971796
case LangAS::opencl_local:
1798-
case LangAS::sycl_local:
17991797
return "__local";
18001798
case LangAS::opencl_private:
1801-
case LangAS::sycl_private:
18021799
return "__private";
18031800
case LangAS::opencl_constant:
1804-
case LangAS::sycl_constant:
18051801
return "__constant";
18061802
case LangAS::opencl_generic:
18071803
return "__generic";

‎clang/lib/Basic/Targets/AMDGPU.cpp

Lines changed: 0 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -48,11 +48,6 @@ const LangASMap AMDGPUTargetInfo::AMDGPUDefIsGenMap = {
4848
Global, // cuda_device
4949
Constant, // cuda_constant
5050
Local, // cuda_shared
51-
Global, // sycl_global
52-
Local, // sycl_local
53-
Constant, // sycl_constant
54-
Private, // sycl_private
55-
Generic, // sycl_generic
5651
Generic, // ptr32_sptr
5752
Generic, // ptr32_uptr
5853
Generic // ptr64
@@ -68,11 +63,6 @@ const LangASMap AMDGPUTargetInfo::AMDGPUDefIsPrivMap = {
6863
Global, // cuda_device
6964
Constant, // cuda_constant
7065
Local, // cuda_shared
71-
Global, // sycl_global
72-
Local, // sycl_local
73-
Constant, // sycl_constant
74-
Private, // sycl_private
75-
Generic, // sycl_generic
7666
Generic, // ptr32_sptr
7767
Generic, // ptr32_uptr
7868
Generic // ptr64

‎clang/lib/Basic/Targets/NVPTX.h

Lines changed: 0 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -33,12 +33,6 @@ static const unsigned NVPTXAddrSpaceMap[] = {
3333
1, // cuda_device
3434
4, // cuda_constant
3535
3, // cuda_shared
36-
1, // sycl_global
37-
3, // sycl_local
38-
4, // sycl_constant
39-
0, // sycl_private
40-
// FIXME: generic has to be added to the target
41-
0, // sycl_generic
4236
0, // ptr32_sptr
4337
0, // ptr32_uptr
4438
0 // ptr64

‎clang/lib/Basic/Targets/SPIR.h

Lines changed: 3 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -33,11 +33,6 @@ static const unsigned SPIRAddrSpaceMap[] = {
3333
0, // cuda_device
3434
0, // cuda_constant
3535
0, // cuda_shared
36-
1, // sycl_global
37-
3, // sycl_local
38-
2, // sycl_constant
39-
0, // sycl_private
40-
4, // sycl_generic
4136
0, // ptr32_sptr
4237
0, // ptr32_uptr
4338
0 // ptr64
@@ -53,11 +48,6 @@ static const unsigned SYCLAddrSpaceMap[] = {
5348
0, // cuda_device
5449
0, // cuda_constant
5550
0, // cuda_shared
56-
1, // sycl_global
57-
3, // sycl_local
58-
2, // sycl_constant
59-
0, // sycl_private
60-
4, // sycl_generic
6151
0, // ptr32_sptr
6252
0, // ptr32_uptr
6353
0 // ptr64
@@ -70,11 +60,9 @@ class LLVM_LIBRARY_VISIBILITY SPIRTargetInfo : public TargetInfo {
7060
TLSSupported = false;
7161
VLASupported = false;
7262
LongWidth = LongAlign = 64;
73-
if (Triple.getEnvironment() == llvm::Triple::SYCLDevice) {
74-
AddrSpaceMap = &SYCLAddrSpaceMap;
75-
} else {
76-
AddrSpaceMap = &SPIRAddrSpaceMap;
77-
}
63+
AddrSpaceMap = (Triple.getEnvironment() == llvm::Triple::SYCLDevice)
64+
? &SYCLAddrSpaceMap
65+
: &SPIRAddrSpaceMap;
7866
UseAddrSpaceMapMangling = true;
7967
HasLegalHalfType = true;
8068
HasFloat16 = true;

‎clang/lib/Basic/Targets/TCE.h

Lines changed: 0 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -40,12 +40,6 @@ static const unsigned TCEOpenCLAddrSpaceMap[] = {
4040
0, // cuda_device
4141
0, // cuda_constant
4242
0, // cuda_shared
43-
3, // sycl_global
44-
4, // sycl_local
45-
5, // sycl_constant
46-
0, // sycl_private
47-
// FIXME: generic has to be added to the target
48-
0, // sycl_generic
4943
0, // ptr32_sptr
5044
0, // ptr32_uptr
5145
0, // ptr64

‎clang/lib/Basic/Targets/X86.h

Lines changed: 0 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -32,11 +32,6 @@ static const unsigned X86AddrSpaceMap[] = {
3232
0, // cuda_device
3333
0, // cuda_constant
3434
0, // cuda_shared
35-
0, // sycl_global
36-
0, // sycl_local
37-
0, // sycl_constant
38-
0, // sycl_private
39-
0, // sycl_generic
4035
270, // ptr32_sptr
4136
271, // ptr32_uptr
4237
272 // ptr64

‎clang/lib/Sema/SemaType.cpp

Lines changed: 3 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -5963,35 +5963,14 @@ static bool BuildAddressSpaceIndex(Sema &S, LangAS &ASIdx,
59635963
llvm::APSInt max(addrSpace.getBitWidth());
59645964
max =
59655965
Qualifiers::MaxAddressSpace - (unsigned)LangAS::FirstTargetAddressSpace;
5966-
59675966
if (addrSpace > max) {
59685967
S.Diag(AttrLoc, diag::err_attribute_address_space_too_high)
59695968
<< (unsigned)max.getZExtValue() << AddrSpace->getSourceRange();
59705969
return false;
59715970
}
59725971

5973-
if (S.LangOpts.SYCLIsDevice && (addrSpace >= 4)) {
5974-
S.Diag(AttrLoc, diag::err_sycl_attribute_address_space_invalid)
5975-
<< AddrSpace->getSourceRange();
5976-
return false;
5977-
}
5978-
5979-
ASIdx = getLangASFromTargetAS(
5980-
static_cast<unsigned>(addrSpace.getZExtValue()));
5981-
5982-
if (S.LangOpts.SYCLIsDevice) {
5983-
ASIdx =
5984-
[](unsigned AS) {
5985-
switch (AS) {
5986-
case 0: return LangAS::sycl_private;
5987-
case 1: return LangAS::sycl_global;
5988-
case 2: return LangAS::sycl_constant;
5989-
case 3: return LangAS::sycl_local;
5990-
case 4: default: llvm_unreachable("Invalid SYCL AS");
5991-
}
5992-
}(static_cast<unsigned>(ASIdx) -
5993-
static_cast<unsigned>(LangAS::FirstTargetAddressSpace));
5994-
}
5972+
ASIdx =
5973+
getLangASFromTargetAS(static_cast<unsigned>(addrSpace.getZExtValue()));
59955974
return true;
59965975
}
59975976

@@ -6117,8 +6096,7 @@ static void HandleAddressSpaceTypeAttribute(QualType &Type,
61176096
Attr.setInvalid();
61186097
} else {
61196098
// The keyword-based type attributes imply which address space to use.
6120-
ASIdx = S.getLangOpts().SYCLIsDevice ?
6121-
Attr.asSYCLLangAS() : Attr.asOpenCLLangAS();
6099+
ASIdx = Attr.asOpenCLLangAS();
61226100
if (ASIdx == LangAS::Default)
61236101
llvm_unreachable("Invalid address space");
61246102

‎clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp

Lines changed: 2 additions & 37 deletions
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@ void foo(int * Data) {}
99
// CHECK-DAG: define spir_func void @[[RAW_PTR:[a-zA-Z0-9_]+]](i32 addrspace(4)* %
1010
void foo2(int * Data) {}
1111
// CHECK-DAG: define spir_func void @[[RAW_PTR2:[a-zA-Z0-9_]+]](i32 addrspace(4)* %
12-
void foo(__attribute__((address_space(3))) int * Data) {}
12+
void foo(__attribute__((opencl_local)) int * Data) {}
1313
// CHECK-DAG: define spir_func void [[LOC_PTR:@[a-zA-Z0-9_]+]](i32 addrspace(3)* %
1414

1515
template<typename T>
@@ -18,12 +18,11 @@ void tmpl(T t){}
1818

1919
void usages() {
2020
// CHECK-DAG: [[GLOB:%[a-zA-Z0-9]+]] = alloca i32 addrspace(1)*
21-
__attribute__((address_space(1))) int *GLOB;
21+
__attribute__((opencl_global)) int *GLOB;
2222
// CHECK-DAG: [[LOC:%[a-zA-Z0-9]+]] = alloca i32 addrspace(3)*
2323
__attribute__((opencl_local)) int *LOC;
2424
// CHECK-DAG: [[NoAS:%[a-zA-Z0-9]+]] = alloca i32 addrspace(4)*
2525
int *NoAS;
26-
2726
// CHECK-DAG: [[PRIV:%[a-zA-Z0-9]+]] = alloca i32*
2827
__attribute__((opencl_private)) int *PRIV;
2928

@@ -94,57 +93,23 @@ void usages() {
9493
// CHECK-DAG: define linkonce_odr spir_func void [[GEN_TMPL]](i32 addrspace(4)* %
9594

9695
void usages2() {
97-
__attribute__((address_space(0))) int *PRIV_NUM;
98-
// CHECK-DAG: [[PRIV_NUM:%[a-zA-Z0-9_]+]] = alloca i32*
99-
__attribute__((address_space(0))) int *PRIV_NUM2;
100-
// CHECK-DAG: [[PRIV_NUM2:%[a-zA-Z0-9_]+]] = alloca i32*
10196
__attribute__((opencl_private)) int *PRIV;
10297
// CHECK-DAG: [[PRIV:%[a-zA-Z0-9_]+]] = alloca i32*
103-
__attribute__((address_space(1))) int *GLOB_NUM;
104-
// CHECK-DAG: [[GLOB_NUM:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(1)*
10598
__attribute__((opencl_global)) int *GLOB;
10699
// CHECK-DAG: [[GLOB:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(1)*
107-
__attribute__((address_space(2))) int *CONST_NUM;
108-
// CHECK-DAG: [[CONST_NUM:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(2)*
109100
__attribute__((opencl_constant)) int *CONST;
110101
// CHECK-DAG: [[CONST:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(2)*
111-
__attribute__((address_space(3))) int *LOCAL_NUM;
112-
// CHECK-DAG: [[LOCAL_NUM:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(3)*
113102
__attribute__((opencl_local)) int *LOCAL;
114103
// CHECK-DAG: [[LOCAL:%[a-zA-Z0-9_]+]] = alloca i32 addrspace(3)*
115104

116-
bar(*PRIV_NUM);
117-
// CHECK-DAG: [[PRIV_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV_NUM]]
118-
// CHECK-DAG: [[PRIV_NUM_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_NUM_LOAD]] to i32 addrspace(4)*
119-
// CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[PRIV_NUM_ASCAST]])
120-
bar(*PRIV_NUM2);
121-
// CHECK-DAG: [[PRIV_NUM2_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV_NUM2]]
122-
// CHECK-DAG: [[PRIV_NUM2_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_NUM2_LOAD]] to i32 addrspace(4)*
123-
// CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[PRIV_NUM2_ASCAST]])
124105
bar(*PRIV);
125106
// CHECK-DAG: [[PRIV_LOAD:%[a-zA-Z0-9]+]] = load i32*, i32** [[PRIV]]
126107
// CHECK-DAG: [[PRIV_ASCAST:%[a-zA-Z0-9]+]] = addrspacecast i32* [[PRIV_LOAD]] to i32 addrspace(4)*
127108
// CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[PRIV_ASCAST]])
128-
bar(*GLOB_NUM);
129-
// CHECK-DAG: [[GLOB_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB_NUM]]
130-
// CHECK-DAG: [[GLOB_NUM_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_NUM_LOAD]] to i32 addrspace(4)*
131-
// CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[GLOB_NUM_CAST]])
132109
bar(*GLOB);
133110
// CHECK-DAG: [[GLOB_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)** [[GLOB]]
134111
// CHECK-DAG: [[GLOB_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(1)* [[GLOB_LOAD]] to i32 addrspace(4)*
135112
// CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[GLOB_CAST]])
136-
bar(*CONST_NUM);
137-
// CHECK-DAG: [[CONST_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(2)*, i32 addrspace(2)** [[CONST_NUM]]
138-
// CHECK-DAG: [[CONST_NUM_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(2)* [[CONST_NUM_LOAD]] to i32 addrspace(4)*
139-
// CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[CONST_NUM_CAST]])
140-
bar(*CONST);
141-
// CHECK-DAG: [[CONST_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(2)*, i32 addrspace(2)** [[CONST]]
142-
// CHECK-DAG: [[CONST_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(2)* [[CONST_LOAD]] to i32 addrspace(4)*
143-
// CHECK-DAG: call spir_func void @[[RAW_REF]](i32 addrspace(4)* dereferenceable(4) [[CONST_CAST]])
144-
bar2(*LOCAL_NUM);
145-
// CHECK-DAG: [[LOCAL_NUM_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOCAL_NUM]]
146-
// CHECK-DAG: [[LOCAL_NUM_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOCAL_NUM_LOAD]] to i32 addrspace(4)*
147-
// CHECK-DAG: call spir_func void @[[RAW_REF2]](i32 addrspace(4)* dereferenceable(4) [[LOCAL_NUM_CAST]])
148113
bar2(*LOCAL);
149114
// CHECK-DAG: [[LOCAL_LOAD:%[a-zA-Z0-9]+]] = load i32 addrspace(3)*, i32 addrspace(3)** [[LOCAL]]
150115
// CHECK-DAG: [[LOCAL_CAST:%[a-zA-Z0-9]+]] = addrspacecast i32 addrspace(3)* [[LOCAL_LOAD]] to i32 addrspace(4)*

‎clang/test/SemaOpenCLCXX/address-space-lambda.cl

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -31,8 +31,8 @@ __kernel void test_qual() {
3131
//CHECK: |-CXXMethodDecl {{.*}} constexpr operator() 'void () const __generic'
3232
auto priv2 = []() __generic {};
3333
priv2();
34-
auto priv3 = []() __global {}; //expected-note{{candidate function not viable: 'this' object is in address space '__private', but method expects object in address space '__global'}} //expected-note{{conversion candidate of type 'void (*)()'}}
35-
priv3(); //expected-error{{no matching function for call to object of type}}
34+
auto priv3 = []() __global {}; //ex pected-note{{candidate function not viable: 'this' object is in address space '__private', but method expects object in address space '__global'}} //ex pected-note{{conversion candidate of type 'void (*)()'}}
35+
priv3(); //ex pected-error{{no matching function for call to object of type}}
3636

3737
__constant auto const1 = []() __private{}; //expected-note{{candidate function not viable: 'this' object is in address space '__constant', but method expects object in address space '__private'}} //expected-note{{conversion candidate of type 'void (*)()'}}
3838
const1(); //expected-error{{no matching function for call to object of type '__constant (lambda at}}

‎clang/test/SemaSYCL/address-space-parameter-conversions.cpp

Lines changed: 1 addition & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -13,7 +13,7 @@ void tmpl(T *t){}
1313
void usages() {
1414
__attribute__((opencl_global)) int *GLOB;
1515
__attribute__((opencl_private)) int *PRIV;
16-
__attribute__((address_space(3))) int *LOC;
16+
__attribute__((opencl_local)) int *LOC;
1717
int *NoAS;
1818

1919
bar(*GLOB);
@@ -53,10 +53,6 @@ void usages() {
5353

5454
// expected-error@+1{{address space is negative}}
5555
__attribute__((address_space(-1))) int *TooLow;
56-
// expected-error@+1{{address space is outside the valid range of values}}
57-
__attribute__((address_space(6))) int *TooHigh;
58-
// expected-error@+1{{address space is outside the valid range of values}}
59-
__attribute__((address_space(4))) int *TriedGeneric;
6056
// expected-error@+1{{unknown type name '__generic'}}
6157
__generic int *IsGeneric;
6258

‎clang/test/SemaTemplate/address_space-dependent.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -43,7 +43,7 @@ void neg() {
4343

4444
template <long int I>
4545
void tooBig() {
46-
__attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388590)}}
46+
__attribute__((address_space(I))) int *bounds; // expected-error {{address space is larger than the maximum supported (8388595)}}
4747
}
4848

4949
template <long int I>

0 commit comments

Comments
 (0)