yaxunl updated this revision to Diff 112288.
yaxunl edited the summary of this revision.
yaxunl added a comment.
Herald added subscribers: nhaehnle, jholewinski.
Add a flag to indicate whether address space qualifier is implicit and only
print explicit address space in diagnostics.
https://reviews.llvm.org/D35082
Files:
include/clang/AST/ASTContext.h
include/clang/AST/Type.h
include/clang/Basic/AddressSpaces.h
lib/AST/ASTContext.cpp
lib/AST/Expr.cpp
lib/AST/ItaniumMangle.cpp
lib/AST/TypePrinter.cpp
lib/Basic/Targets/AMDGPU.cpp
lib/Basic/Targets/NVPTX.h
lib/Basic/Targets/SPIR.h
lib/Basic/Targets/TCE.h
lib/CodeGen/CGDecl.cpp
lib/Sema/SemaChecking.cpp
lib/Sema/SemaDecl.cpp
lib/Sema/SemaType.cpp
test/CodeGenOpenCL/address-spaces-mangling.cl
test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
test/SemaOpenCL/address-spaces.cl
test/SemaOpenCL/atomic-ops.cl
test/SemaOpenCL/cl20-device-side-enqueue.cl
test/SemaOpenCL/invalid-block.cl
test/SemaOpenCL/invalid-pipes-cl2.0.cl
test/SemaOpenCL/null_literal.cl
Index: test/SemaOpenCL/null_literal.cl
===================================================================
--- test/SemaOpenCL/null_literal.cl
+++ test/SemaOpenCL/null_literal.cl
@@ -1,5 +1,5 @@
// RUN: %clang_cc1 -verify %s
-// RUN: %clang_cc1 -cl-std=CL2.0 -DCL20 -verify %s
+// RUN: %clang_cc1 -cl-std=CL2.0 -verify %s
#define NULL ((void*)0)
@@ -13,9 +13,17 @@
constant int* ptr4 = (global void*)0; // expected-error{{initializing '__constant int *' with an expression of type '__global void *' changes address space of pointer}}
-#ifdef CL20
+#if __OPENCL_C_VERSION__ >= 200
+
// Accept explicitly pointer to generic address space in OpenCL v2.0.
global int* ptr5 = (generic void*)0;
+
+global int *ptr = (private void *)0; // expected-error{{initializing '__global int *' with an expression of type '__private void *' changes address space of pointer}}
+
+#else
+
+global int *ptr = (private void *)0;
+
#endif
global int* ptr6 = (local void*)0; // expected-error{{initializing '__global int *' with an expression of type '__local void *' changes address space of pointer}}
Index: test/SemaOpenCL/invalid-pipes-cl2.0.cl
===================================================================
--- test/SemaOpenCL/invalid-pipes-cl2.0.cl
+++ test/SemaOpenCL/invalid-pipes-cl2.0.cl
@@ -3,7 +3,7 @@
global pipe int gp; // expected-error {{type '__global read_only pipe int' can only be used as a function parameter in OpenCL}}
global reserve_id_t rid; // expected-error {{the '__global reserve_id_t' type cannot be used to declare a program scope variable}}
-extern pipe write_only int get_pipe(); // expected-error {{type '__global write_only pipe int ()' can only be used as a function parameter in OpenCL}}
+extern pipe write_only int get_pipe(); // expected-error {{type 'write_only pipe int ()' can only be used as a function parameter in OpenCL}}
kernel void test_invalid_reserved_id(reserve_id_t ID) { // expected-error {{'reserve_id_t' cannot be used as the type of a kernel parameter}}
}
Index: test/SemaOpenCL/invalid-block.cl
===================================================================
--- test/SemaOpenCL/invalid-block.cl
+++ test/SemaOpenCL/invalid-block.cl
@@ -12,7 +12,7 @@
};
f0(bl1);
f0(bl2);
- bl1 = bl2; // expected-error{{invalid operands to binary expression ('int (__generic ^const)(void)' and 'int (__generic ^const)(void)')}}
+ bl1 = bl2; // expected-error{{invalid operands to binary expression ('int (^const)(void)' and 'int (^const)(void)')}}
int (^const bl3)(); // expected-error{{invalid block variable declaration - must be initialized}}
}
@@ -28,10 +28,10 @@
// A block cannot be the return value of a function.
typedef int (^bl_t)(void);
-bl_t f3(bl_t bl); // expected-error{{declaring function return value of type 'bl_t' (aka 'int (__generic ^const)(void)') is not allowed}}
+bl_t f3(bl_t bl); // expected-error{{declaring function return value of type 'bl_t' (aka 'int (^const)(void)') is not allowed}}
struct bl_s {
- int (^bl)(void); // expected-error {{the 'int (__generic ^const)(void)' type cannot be used to declare a structure or union field}}
+ int (^bl)(void); // expected-error {{the 'int (^const)(void)' type cannot be used to declare a structure or union field}}
};
void f4() {
@@ -53,18 +53,18 @@
bl2_t bl2 = ^(int i) {
return 2;
};
- bl2_t arr[] = {bl1, bl2}; // expected-error {{array of 'bl2_t' (aka 'int (__generic ^const)(int)') type is invalid in OpenCL}}
+ bl2_t arr[] = {bl1, bl2}; // expected-error {{array of 'bl2_t' (aka 'int (^const)(int)') type is invalid in OpenCL}}
int tmp = i ? bl1(i) // expected-error {{block type cannot be used as expression in ternary expression in OpenCL}}
: bl2(i); // expected-error {{block type cannot be used as expression in ternary expression in OpenCL}}
}
// A block pointer type and all pointer operations are disallowed
-void f6(bl2_t *bl_ptr) { // expected-error{{pointer to type '__generic bl2_t' (aka 'int (__generic ^const __generic)(int)') is invalid in OpenCL}}
+void f6(bl2_t *bl_ptr) { // expected-error{{pointer to type 'bl2_t' (aka 'int (^const)(int)') is invalid in OpenCL}}
bl2_t bl = ^(int i) {
return 1;
};
- bl2_t *p; // expected-error {{pointer to type '__generic bl2_t' (aka 'int (__generic ^const __generic)(int)') is invalid in OpenCL}}
- *bl; // expected-error {{invalid argument type 'bl2_t' (aka 'int (__generic ^const)(int)') to unary expression}}
- &bl; // expected-error {{invalid argument type 'bl2_t' (aka 'int (__generic ^const)(int)') to unary expression}}
+ bl2_t *p; // expected-error {{pointer to type 'bl2_t' (aka 'int (^const)(int)') is invalid in OpenCL}}
+ *bl; // expected-error {{invalid argument type 'bl2_t' (aka 'int (^const)(int)') to unary expression}}
+ &bl; // expected-error {{invalid argument type 'bl2_t' (aka 'int (^const)(int)') to unary expression}}
}
// A block can't reference another block
kernel void f7() {
Index: test/SemaOpenCL/cl20-device-side-enqueue.cl
===================================================================
--- test/SemaOpenCL/cl20-device-side-enqueue.cl
+++ test/SemaOpenCL/cl20-device-side-enqueue.cl
@@ -222,7 +222,7 @@
kernel void bar(global int *buf)
{
- ndrange_t n;
+ __private ndrange_t n;
buf[0] = get_kernel_sub_group_count_for_ndrange(n, ^(){});
buf[0] = get_kernel_sub_group_count_for_ndrange(0, ^(){}); // expected-error{{illegal call to 'get_kernel_sub_group_count_for_ndrange', expected 'ndrange_t' argument type}}
buf[0] = get_kernel_sub_group_count_for_ndrange(n, 1); // expected-error{{illegal call to 'get_kernel_sub_group_count_for_ndrange', expected block argument type}}
Index: test/SemaOpenCL/atomic-ops.cl
===================================================================
--- test/SemaOpenCL/atomic-ops.cl
+++ test/SemaOpenCL/atomic-ops.cl
@@ -41,24 +41,24 @@
intptr_t *P, float *D, struct S *s1, struct S *s2,
global atomic_int *i_g, local atomic_int *i_l, private atomic_int *i_p,
constant atomic_int *i_c) {
- __opencl_atomic_init(I, 5); // expected-error {{address argument to atomic operation must be a pointer to _Atomic type ('__generic int *' invalid)}}
- __opencl_atomic_init(ci, 5); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
+ __opencl_atomic_init(I, 5); // expected-error {{address argument to atomic operation must be a pointer to _Atomic type ('int *' invalid)}}
+ __opencl_atomic_init(ci, 5); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const atomic_int *' (aka 'const _Atomic(int) *') invalid)}}
__opencl_atomic_load(0); // expected-error {{too few arguments to function call, expected 3, have 1}}
__opencl_atomic_load(0, 0, 0, 0); // expected-error {{too many arguments to function call, expected 3, have 4}}
__opencl_atomic_store(0,0,0,0); // expected-error {{address argument to atomic builtin must be a pointer}}
- __opencl_atomic_store((int *)0, 0, 0, 0); // expected-error {{address argument to atomic operation must be a pointer to _Atomic type ('__generic int *' invalid)}}
+ __opencl_atomic_store((int *)0, 0, 0, 0); // expected-error {{address argument to atomic operation must be a pointer to _Atomic type ('int *' invalid)}}
__opencl_atomic_store(i, 0, memory_order_relaxed, memory_scope_work_group);
- __opencl_atomic_store(ci, 0, memory_order_relaxed, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
+ __opencl_atomic_store(ci, 0, memory_order_relaxed, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const atomic_int *' (aka 'const _Atomic(int) *') invalid)}}
__opencl_atomic_store(i_g, 0, memory_order_relaxed, memory_scope_work_group);
__opencl_atomic_store(i_l, 0, memory_order_relaxed, memory_scope_work_group);
__opencl_atomic_store(i_p, 0, memory_order_relaxed, memory_scope_work_group);
__opencl_atomic_store(i_c, 0, memory_order_relaxed, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-constant _Atomic type ('__constant atomic_int *' (aka '__constant _Atomic(int) *') invalid)}}
__opencl_atomic_load(i, memory_order_seq_cst, memory_scope_work_group);
__opencl_atomic_load(p, memory_order_seq_cst, memory_scope_work_group);
__opencl_atomic_load(d, memory_order_seq_cst, memory_scope_work_group);
- __opencl_atomic_load(ci, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
+ __opencl_atomic_load(ci, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const atomic_int *' (aka 'const _Atomic(int) *') invalid)}}
__opencl_atomic_store(i, 1, memory_order_seq_cst, memory_scope_work_group);
__opencl_atomic_store(p, 1, memory_order_seq_cst, memory_scope_work_group);
@@ -69,35 +69,35 @@
__opencl_atomic_fetch_add(i, 1, memory_order_seq_cst, memory_scope_work_group);
__opencl_atomic_fetch_add(p, 1, memory_order_seq_cst, memory_scope_work_group);
- __opencl_atomic_fetch_add(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer or pointer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
+ __opencl_atomic_fetch_add(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer or pointer ('atomic_float *' (aka '_Atomic(float) *') invalid)}}
__opencl_atomic_fetch_and(i, 1, memory_order_seq_cst, memory_scope_work_group);
__opencl_atomic_fetch_and(p, 1, memory_order_seq_cst, memory_scope_work_group);
- __opencl_atomic_fetch_and(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to bitwise atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
+ __opencl_atomic_fetch_and(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to bitwise atomic operation must be a pointer to atomic integer ('atomic_float *' (aka '_Atomic(float) *') invalid)}}
__opencl_atomic_fetch_min(i, 1, memory_order_seq_cst, memory_scope_work_group);
__opencl_atomic_fetch_max(i, 1, memory_order_seq_cst, memory_scope_work_group);
- __opencl_atomic_fetch_min(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer or pointer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
- __opencl_atomic_fetch_max(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer or pointer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
+ __opencl_atomic_fetch_min(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer or pointer ('atomic_float *' (aka '_Atomic(float) *') invalid)}}
+ __opencl_atomic_fetch_max(d, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer or pointer ('atomic_float *' (aka '_Atomic(float) *') invalid)}}
bool cmpexch_1 = __opencl_atomic_compare_exchange_strong(i, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
bool cmpexch_2 = __opencl_atomic_compare_exchange_strong(p, P, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
- bool cmpexch_3 = __opencl_atomic_compare_exchange_strong(d, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *' to parameter of type '__generic float *'}}
- (void)__opencl_atomic_compare_exchange_strong(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const __generic int *' to parameter of type '__generic int *' discards qualifiers}}
+ bool cmpexch_3 = __opencl_atomic_compare_exchange_strong(d, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing 'int *' to parameter of type '__generic float *'}}
+ (void)__opencl_atomic_compare_exchange_strong(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const int *' to parameter of type '__generic int *' discards qualifiers}}
bool cmpexchw_1 = __opencl_atomic_compare_exchange_weak(i, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
bool cmpexchw_2 = __opencl_atomic_compare_exchange_weak(p, P, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
- bool cmpexchw_3 = __opencl_atomic_compare_exchange_weak(d, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *' to parameter of type '__generic float *'}}
- (void)__opencl_atomic_compare_exchange_weak(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const __generic int *' to parameter of type '__generic int *' discards qualifiers}}
+ bool cmpexchw_3 = __opencl_atomic_compare_exchange_weak(d, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing 'int *' to parameter of type '__generic float *'}}
+ (void)__opencl_atomic_compare_exchange_weak(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const int *' to parameter of type '__generic int *' discards qualifiers}}
// Pointers to different address spaces are allowed.
bool cmpexch_10 = __opencl_atomic_compare_exchange_strong((global atomic_int *)0x308, (constant int *)0x309, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
- __opencl_atomic_init(ci, 0); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
- __opencl_atomic_store(ci, 0, memory_order_release, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
- __opencl_atomic_load(ci, memory_order_acquire, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
+ __opencl_atomic_init(ci, 0); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const atomic_int *' (aka 'const _Atomic(int) *') invalid)}}
+ __opencl_atomic_store(ci, 0, memory_order_release, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const atomic_int *' (aka 'const _Atomic(int) *') invalid)}}
+ __opencl_atomic_load(ci, memory_order_acquire, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const atomic_int *' (aka 'const _Atomic(int) *') invalid)}}
__opencl_atomic_init(&gn, 456);
- __opencl_atomic_init(&gn, (void*)0); // expected-warning{{incompatible pointer to integer conversion passing '__generic void *' to parameter of type 'int'}}
+ __opencl_atomic_init(&gn, (void*)0); // expected-warning{{incompatible pointer to integer conversion passing 'void *' to parameter of type 'int'}}
}
void memory_checks(atomic_int *Ap, int *p, int val) {
Index: test/SemaOpenCL/address-spaces.cl
===================================================================
--- test/SemaOpenCL/address-spaces.cl
+++ test/SemaOpenCL/address-spaces.cl
@@ -1,42 +1,65 @@
// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -pedantic -fsyntax-only
__constant int ci = 1;
__kernel void foo(__global int *gip) {
__local int li;
__local int lj = 2; // expected-error {{'__local' variable cannot have an initializer}}
int *ip;
+#if __OPENCL_C_VERSION__ < 200
ip = gip; // expected-error {{assigning '__global int *' to 'int *' changes address space of pointer}}
ip = &li; // expected-error {{assigning '__local int *' to 'int *' changes address space of pointer}}
ip = &ci; // expected-error {{assigning '__constant int *' to 'int *' changes address space of pointer}}
+#else
+ ip = gip;
+ ip = &li;
+ ip = &ci; // expected-error {{assigning '__constant int *' to 'int *' changes address space of pointer}}
+#endif
}
void explicit_cast(global int* g, local int* l, constant int* c, private int* p, const constant int *cc)
{
g = (global int*) l; // expected-error {{casting '__local int *' to type '__global int *' changes address space of pointer}}
g = (global int*) c; // expected-error {{casting '__constant int *' to type '__global int *' changes address space of pointer}}
g = (global int*) cc; // expected-error {{casting 'const __constant int *' to type '__global int *' changes address space of pointer}}
- g = (global int*) p; // expected-error {{casting 'int *' to type '__global int *' changes address space of pointer}}
+ g = (global int*) p; // expected-error {{casting '__private int *' to type '__global int *' changes address space of pointer}}
l = (local int*) g; // expected-error {{casting '__global int *' to type '__local int *' changes address space of pointer}}
l = (local int*) c; // expected-error {{casting '__constant int *' to type '__local int *' changes address space of pointer}}
l = (local int*) cc; // expected-error {{casting 'const __constant int *' to type '__local int *' changes address space of pointer}}
- l = (local int*) p; // expected-error {{casting 'int *' to type '__local int *' changes address space of pointer}}
+ l = (local int*) p; // expected-error {{casting '__private int *' to type '__local int *' changes address space of pointer}}
c = (constant int*) g; // expected-error {{casting '__global int *' to type '__constant int *' changes address space of pointer}}
c = (constant int*) l; // expected-error {{casting '__local int *' to type '__constant int *' changes address space of pointer}}
- c = (constant int*) p; // expected-error {{casting 'int *' to type '__constant int *' changes address space of pointer}}
+ c = (constant int*) p; // expected-error {{casting '__private int *' to type '__constant int *' changes address space of pointer}}
- p = (private int*) g; // expected-error {{casting '__global int *' to type 'int *' changes address space of pointer}}
- p = (private int*) l; // expected-error {{casting '__local int *' to type 'int *' changes address space of pointer}}
- p = (private int*) c; // expected-error {{casting '__constant int *' to type 'int *' changes address space of pointer}}
- p = (private int*) cc; // expected-error {{casting 'const __constant int *' to type 'int *' changes address space of pointer}}
+ p = (private int*) g; // expected-error {{casting '__global int *' to type '__private int *' changes address space of pointer}}
+ p = (private int*) l; // expected-error {{casting '__local int *' to type '__private int *' changes address space of pointer}}
+ p = (private int*) c; // expected-error {{casting '__constant int *' to type '__private int *' changes address space of pointer}}
+ p = (private int*) cc; // expected-error {{casting 'const __constant int *' to type '__private int *' changes address space of pointer}}
}
void ok_explicit_casts(global int *g, global int* g2, local int* l, local int* l2, private int* p, private int* p2)
{
g = (global int*) g2;
l = (local int*) l2;
p = (private int*) p2;
}
+
+__private int func_return_priv(void); //expected-error {{return value cannot be qualified with address space}}
+__global int func_return_global(void); //expected-error {{return value cannot be qualified with address space}}
+__local int func_return_local(void); //expected-error {{return value cannot be qualified with address space}}
+__constant int func_return_constant(void); //expected-error {{return value cannot be qualified with address space}}
+#if __OPENCL_C_VERSION__ >= 200
+__generic int func_return_generic(void); //expected-error {{return value cannot be qualified with address space}}
+#endif
+
+void func_multiple_addr(void) {
+ typedef __private int private_int_t;
+ __local __private int var1; // expected-error {{multiple address spaces specified for type}}
+ __local __private int *var2; // expected-error {{multiple address spaces specified for type}}
+ __local private_int_t var3; // expected-error {{multiple address spaces specified for type}}
+ __local private_int_t *var4; // expected-error {{multiple address spaces specified for type}}
+}
Index: test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
===================================================================
--- test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
+++ test/SemaOpenCL/address-spaces-conversions-cl2.0.cl
@@ -76,7 +76,7 @@
AS int *var_init4 = arg_priv;
#ifndef GENERIC
-// expected-error-re@-2{{initializing '__{{global|constant}} int *' with an expression of type 'int *' changes address space of pointer}}
+// expected-error-re@-2{{initializing '__{{global|constant}} int *' with an expression of type '__private int *' changes address space of pointer}}
#endif
AS int *var_init5 = arg_gen;
@@ -101,7 +101,7 @@
AS int *var_cast4 = (AS int *)arg_priv;
#ifndef GENERIC
-// expected-error-re@-2{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re@-2{{casting '__private int *' to type '__{{global|constant}} int *' changes address space of pointer}}
#endif
AS int *var_cast5 = (AS int *)arg_gen;
@@ -127,7 +127,7 @@
var_impl = arg_priv;
#ifndef GENERIC
-// expected-error-re@-2{{assigning 'int *' to '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re@-2{{assigning '__private int *' to '__{{global|constant}} int *' changes address space of pointer}}
#endif
var_impl = arg_gen;
@@ -152,7 +152,7 @@
var_cast4 = (AS int *)arg_priv;
#ifndef GENERIC
-// expected-error-re@-2{{casting 'int *' to type '__{{global|constant}} int *' changes address space of pointer}}
+// expected-error-re@-2{{casting '__private int *' to type '__{{global|constant}} int *' changes address space of pointer}}
#endif
var_cast5 = (AS int *)arg_gen;
@@ -178,7 +178,7 @@
b = var_cmp <= arg_priv;
#ifndef GENERIC
-// expected-error-re@-2{{comparison between ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}}
+// expected-error-re@-2{{comparison between ('__{{global|constant}} int *' and '__private int *') which are pointers to non-overlapping address spaces}}
#endif
b = var_cmp >= arg_gen;
@@ -204,7 +204,7 @@
b = var_sub - arg_priv;
#ifndef GENERIC
-// expected-error-re@-2{{arithmetic operation with operands of type ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}}
+// expected-error-re@-2{{arithmetic operation with operands of type ('__{{global|constant}} int *' and '__private int *') which are pointers to non-overlapping address spaces}}
#endif
b = var_sub - arg_gen;
@@ -224,7 +224,7 @@
// expected-error-re@-2{{passing '__{{global|generic}} int *' to parameter of type '__constant int *' changes address space of pointer}}
#endif
- f_priv(var_sub); // expected-error-re{{passing '__{{global|constant|generic}} int *' to parameter of type 'int *' changes address space of pointer}}
+ f_priv(var_sub); // expected-error-re{{passing '__{{global|constant|generic}} int *' to parameter of type '__private int *' changes address space of pointer}}
f_gen(var_sub);
#ifdef CONSTANT
@@ -256,7 +256,7 @@
private int *var_priv;
var_gen = 0 ? var_cond : var_priv;
#ifndef GENERIC
-// expected-error-re@-2{{conditional operator with the second and third operands of type ('__{{global|constant}} int *' and 'int *') which are pointers to non-overlapping address spaces}}
+// expected-error-re@-2{{conditional operator with the second and third operands of type ('__{{global|constant}} int *' and '__private int *') which are pointers to non-overlapping address spaces}}
#endif
var_gen = 0 ? var_cond : var_gen;
@@ -293,9 +293,9 @@
private char *var_priv_ch;
var_void_gen = 0 ? var_cond : var_priv_ch;
#ifndef GENERIC
-// expected-error-re@-2{{conditional operator with the second and third operands of type ('__{{global|constant}} int *' and 'char *') which are pointers to non-overlapping address spaces}}
+// expected-error-re@-2{{conditional operator with the second and third operands of type ('__{{global|constant}} int *' and '__private char *') which are pointers to non-overlapping address spaces}}
#else
-// expected-warning@-4{{pointer type mismatch ('__generic int *' and 'char *')}}
+// expected-warning@-4{{pointer type mismatch ('__generic int *' and '__private char *')}}
#endif
generic char *var_gen_ch;
Index: test/CodeGenOpenCL/address-spaces-mangling.cl
===================================================================
--- test/CodeGenOpenCL/address-spaces-mangling.cl
+++ test/CodeGenOpenCL/address-spaces-mangling.cl
@@ -1,5 +1,7 @@
-// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=yes -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefix=ASMANG %s
-// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=no -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefix=NOASMANG %s
+// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=yes -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefixes=ASMANG,ASMAN10 %s
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -ffake-address-space-map -faddress-space-map-mangling=yes -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefixes=ASMANG,ASMAN20 %s
+// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=no -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefixes=NOASMANG,NOASMAN10 %s
+// RUN: %clang_cc1 %s -cl-std=CL2.0 -ffake-address-space-map -faddress-space-map-mangling=no -triple %itanium_abi_triple -emit-llvm -o - | FileCheck -check-prefixes=NOASMANG,NOASMAN20 %s
// We check that the address spaces are mangled the same in both version of OpenCL
// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=CL2.0 -emit-llvm -o - | FileCheck -check-prefix=OCL-20 %s
@@ -10,15 +12,17 @@
// warnings, but we do want it for comparison purposes.
__attribute__((overloadable))
void ff(int *arg) { }
-// ASMANG: @_Z2ffPi
-// NOASMANG: @_Z2ffPi
+// ASMANG10: @_Z2ffPi
+// ASMANG20: @_Z2ffPU3AS4i
+// NOASMANG10: @_Z2ffPi
+// NOASMANG20: @_Z2ffPU9CLgenerici
// OCL-20-DAG: @_Z2ffPU3AS4i
// OCL-12-DAG: @_Z2ffPi
__attribute__((overloadable))
void f(private int *arg) { }
// ASMANG: @_Z1fPi
-// NOASMANG: @_Z1fPi
+// NOASMANG: @_Z1fPU9CLprivatei
// OCL-20-DAG: @_Z1fPi
// OCL-12-DAG: @_Z1fPi
@@ -42,3 +46,11 @@
// NOASMANG: @_Z1fPU10CLconstanti
// OCL-20-DAG: @_Z1fPU3AS2i
// OCL-12-DAG: @_Z1fPU3AS2i
+
+#if __OPENCL_C_VERSION__ >= 200
+__attribute__((overloadable))
+void f(generic int *arg) { }
+// ASMANG20: @_Z1fPU3AS4i
+// NOASMANG20: @_Z1fPU9CLgenerici
+// OCL-20-DAG: @_Z1fPU3AS4i
+#endif
Index: lib/Sema/SemaType.cpp
===================================================================
--- lib/Sema/SemaType.cpp
+++ lib/Sema/SemaType.cpp
@@ -5568,7 +5568,7 @@
ASIdx = LangAS::opencl_generic; break;
default:
assert(Attr.getKind() == AttributeList::AT_OpenCLPrivateAddressSpace);
- ASIdx = 0; break;
+ ASIdx = LangAS::opencl_private; break;
}
}
@@ -6971,6 +6971,38 @@
}
}
+ if (state.getSema().getLangOpts().OpenCL && !hasOpenCLAddressSpace &&
+ type.getAddressSpace() == LangAS::Default &&
+ (TAL == TAL_DeclSpec || TAL == TAL_DeclChunk)) {
+ Declarator &D = state.getDeclarator();
+ // Put OpenCL automatic variable in private address space.
+ if (state.getCurrentChunkIndex() == 0 && !D.isFunctionDeclarator() &&
+ !D.isFunctionDefinition() &&
+ D.getDeclSpec().getStorageClassSpec() != DeclSpec::SCS_typedef &&
+ !type->isVoidType()) {
+ if (D.getContext() == Declarator::BlockContext ||
+ D.getContext() == Declarator::ForContext ||
+ D.getContext() == Declarator::InitStmtContext ||
+ D.getContext() == Declarator::ConditionContext ||
+ D.getContext() == Declarator::PrototypeContext) {
+ if (D.getDeclSpec().getStorageClassSpec() != DeclSpec::SCS_static) {
+ type = state.getSema().Context.getAddrSpaceQualType(
+ type, LangAS::opencl_private, true);
+ }
+ }
+ // OpenCL v1.2 s6.5:
+ // The generic address space name for arguments to a function in a
+ // program, or local variables of a function is __private. All function
+ // arguments shall be in the __private address space.
+ } else if (state.getCurrentChunkIndex() > 0 &&
+ state.getSema().getLangOpts().OpenCLVersion <= 120 &&
+ (D.getTypeObject(state.getCurrentChunkIndex() - 1).Kind ==
+ DeclaratorChunk::Pointer)) {
+ type = state.getSema().Context.getAddrSpaceQualType(
+ type, LangAS::opencl_private, true);
+ }
+ }
+
// If address space is not set, OpenCL 2.0 defines non private default
// address spaces for some cases:
// OpenCL 2.0, section 6.5:
@@ -6981,28 +7013,37 @@
// Pointers that are declared without pointing to a named address space point
// to the generic address space.
if (state.getSema().getLangOpts().OpenCLVersion >= 200 &&
- !hasOpenCLAddressSpace && type.getAddressSpace() == 0 &&
+ !hasOpenCLAddressSpace && type.getAddressSpace() == LangAS::Default &&
(TAL == TAL_DeclSpec || TAL == TAL_DeclChunk)) {
Declarator &D = state.getDeclarator();
if (state.getCurrentChunkIndex() > 0 &&
(D.getTypeObject(state.getCurrentChunkIndex() - 1).Kind ==
DeclaratorChunk::Pointer ||
D.getTypeObject(state.getCurrentChunkIndex() - 1).Kind ==
DeclaratorChunk::BlockPointer)) {
type = state.getSema().Context.getAddrSpaceQualType(
- type, LangAS::opencl_generic);
+ type, LangAS::opencl_generic, true);
} else if (state.getCurrentChunkIndex() == 0 &&
- D.getContext() == Declarator::FileContext &&
!D.isFunctionDeclarator() && !D.isFunctionDefinition() &&
D.getDeclSpec().getStorageClassSpec() != DeclSpec::SCS_typedef &&
- !type->isSamplerT())
- type = state.getSema().Context.getAddrSpaceQualType(
- type, LangAS::opencl_global);
- else if (state.getCurrentChunkIndex() == 0 &&
- D.getContext() == Declarator::BlockContext &&
- D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_static)
+ !type->isSamplerT()) {
+ if (D.getContext() == Declarator::FileContext) {
+ type = state.getSema().Context.getAddrSpaceQualType(
+ type, LangAS::opencl_global, true);
+ } else if (D.getContext() == Declarator::BlockContext) {
+ if (D.getDeclSpec().getStorageClassSpec() != DeclSpec::SCS_static) {
+ type = state.getSema().Context.getAddrSpaceQualType(
+ type, LangAS::opencl_private, true);
+ } else {
+ type = state.getSema().Context.getAddrSpaceQualType(
+ type, LangAS::opencl_global, true);
+ }
+ }
+ } else if (state.getCurrentChunkIndex() == 0 &&
+ D.getContext() == Declarator::BlockContext &&
+ D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_static)
type = state.getSema().Context.getAddrSpaceQualType(
- type, LangAS::opencl_global);
+ type, LangAS::opencl_global, true);
}
}
Index: lib/Sema/SemaDecl.cpp
===================================================================
--- lib/Sema/SemaDecl.cpp
+++ lib/Sema/SemaDecl.cpp
@@ -6237,7 +6237,7 @@
// The event type cannot be used with the __local, __constant and __global
// address space qualifiers.
if (R->isEventT()) {
- if (R.getAddressSpace()) {
+ if (R.getAddressSpace() != LangAS::opencl_private) {
Diag(D.getLocStart(), diag::err_event_t_addr_space_qual);
D.setInvalidType();
}
@@ -7340,7 +7340,7 @@
return;
}
}
- } else if (T.getAddressSpace() != LangAS::Default) {
+ } else if (T.getAddressSpace() != LangAS::opencl_private) {
// Do not allow other address spaces on automatic variable.
Diag(NewVD->getLocation(), diag::err_as_qualified_auto_decl) << 1;
NewVD->setInvalidDecl();
@@ -7975,7 +7975,8 @@
if (PointeeType->isPointerType())
return PtrPtrKernelParam;
if (PointeeType.getAddressSpace() == LangAS::opencl_generic ||
- PointeeType.getAddressSpace() == 0)
+ PointeeType.getAddressSpace() == LangAS::opencl_private ||
+ PointeeType.getAddressSpace() == LangAS::Default)
return InvalidAddrSpacePtrKernelParam;
return PtrKernelParam;
}
@@ -8745,9 +8746,7 @@
// OpenCL v1.1 s6.5: Using an address space qualifier in a function return
// type declaration will generate a compilation error.
unsigned AddressSpace = NewFD->getReturnType().getAddressSpace();
- if (AddressSpace == LangAS::opencl_local ||
- AddressSpace == LangAS::opencl_global ||
- AddressSpace == LangAS::opencl_constant) {
+ if (AddressSpace != LangAS::Default) {
Diag(NewFD->getLocation(),
diag::err_opencl_return_value_with_address_space);
NewFD->setInvalidDecl();
@@ -11857,13 +11856,13 @@
// duration shall not be qualified by an address-space qualifier."
// Since all parameters have automatic store duration, they can not have
// an address space.
- if (T.getAddressSpace() != 0) {
- // OpenCL allows function arguments declared to be an array of a type
- // to be qualified with an address space.
- if (!(getLangOpts().OpenCL && T->isArrayType())) {
- Diag(NameLoc, diag::err_arg_with_address_space);
- New->setInvalidDecl();
- }
+ if (T.getAddressSpace() != LangAS::Default &&
+ // OpenCL allows function arguments declared to be an array of a type
+ // to be qualified with an address space.
+ !(getLangOpts().OpenCL &&
+ (T->isArrayType() || T.getAddressSpace() == LangAS::opencl_private))) {
+ Diag(NameLoc, diag::err_arg_with_address_space);
+ New->setInvalidDecl();
}
return New;
Index: lib/Sema/SemaChecking.cpp
===================================================================
--- lib/Sema/SemaChecking.cpp
+++ lib/Sema/SemaChecking.cpp
@@ -318,7 +318,7 @@
// First argument is an ndrange_t type.
Expr *NDRangeArg = TheCall->getArg(0);
- if (NDRangeArg->getType().getAsString() != "ndrange_t") {
+ if (NDRangeArg->getType().getUnqualifiedType().getAsString() != "ndrange_t") {
S.Diag(NDRangeArg->getLocStart(),
diag::err_opencl_builtin_expected_type)
<< TheCall->getDirectCallee() << "'ndrange_t'";
@@ -762,8 +762,11 @@
case Builtin::BIto_local:
Qual.setAddressSpace(LangAS::opencl_local);
break;
+ case Builtin::BIto_private:
+ Qual.setAddressSpace(LangAS::opencl_private);
+ break;
default:
- Qual.removeAddressSpace();
+ llvm_unreachable("Invalid builtin function");
}
Call->setType(S.Context.getPointerType(S.Context.getQualifiedType(
RT.getUnqualifiedType(), Qual)));
Index: lib/CodeGen/CGDecl.cpp
===================================================================
--- lib/CodeGen/CGDecl.cpp
+++ lib/CodeGen/CGDecl.cpp
@@ -956,7 +956,9 @@
CodeGenFunction::AutoVarEmission
CodeGenFunction::EmitAutoVarAlloca(const VarDecl &D) {
QualType Ty = D.getType();
- assert(Ty.getAddressSpace() == LangAS::Default);
+ assert(
+ Ty.getAddressSpace() == LangAS::Default ||
+ (Ty.getAddressSpace() == LangAS::opencl_private && getLangOpts().OpenCL));
AutoVarEmission emission(D);
Index: lib/Basic/Targets/TCE.h
===================================================================
--- lib/Basic/Targets/TCE.h
+++ lib/Basic/Targets/TCE.h
@@ -35,6 +35,7 @@
3, // opencl_global
4, // opencl_local
5, // opencl_constant
+ 0, // opencl_private
// FIXME: generic has to be added to the target
0, // opencl_generic
0, // cuda_device
Index: lib/Basic/Targets/SPIR.h
===================================================================
--- lib/Basic/Targets/SPIR.h
+++ lib/Basic/Targets/SPIR.h
@@ -27,6 +27,7 @@
1, // opencl_global
3, // opencl_local
2, // opencl_constant
+ 0, // opencl_private
4, // opencl_generic
0, // cuda_device
0, // cuda_constant
Index: lib/Basic/Targets/NVPTX.h
===================================================================
--- lib/Basic/Targets/NVPTX.h
+++ lib/Basic/Targets/NVPTX.h
@@ -28,6 +28,7 @@
1, // opencl_global
3, // opencl_local
4, // opencl_constant
+ 0, // opencl_private
// FIXME: generic has to be added to the target
0, // opencl_generic
1, // cuda_device
Index: lib/Basic/Targets/AMDGPU.cpp
===================================================================
--- lib/Basic/Targets/AMDGPU.cpp
+++ lib/Basic/Targets/AMDGPU.cpp
@@ -47,6 +47,7 @@
1, // opencl_global
3, // opencl_local
2, // opencl_constant
+ 0, // opencl_private
4, // opencl_generic
1, // cuda_device
2, // cuda_constant
@@ -58,6 +59,7 @@
1, // opencl_global
3, // opencl_local
2, // opencl_constant
+ 5, // opencl_private
0, // opencl_generic
1, // cuda_device
2, // cuda_constant
@@ -69,6 +71,7 @@
1, // opencl_global
3, // opencl_local
2, // opencl_constant
+ 0, // opencl_private
4, // opencl_generic
1, // cuda_device
2, // cuda_constant
@@ -80,6 +83,7 @@
1, // opencl_global
3, // opencl_local
2, // opencl_constant
+ 5, // opencl_private
0, // opencl_generic
1, // cuda_device
2, // cuda_constant
Index: lib/AST/TypePrinter.cpp
===================================================================
--- lib/AST/TypePrinter.cpp
+++ lib/AST/TypePrinter.cpp
@@ -1660,11 +1660,12 @@
OS << "__unaligned";
addSpace = true;
}
- if (unsigned addrspace = getAddressSpace()) {
- if (addSpace)
- OS << ' ';
- addSpace = true;
- switch (addrspace) {
+ if (!getImplicitAddressSpaceFlag()) {
+ if (unsigned addrspace = getAddressSpace()) {
+ if (addSpace)
+ OS << ' ';
+ addSpace = true;
+ switch (addrspace) {
case LangAS::opencl_global:
OS << "__global";
break;
@@ -1675,6 +1676,9 @@
case LangAS::cuda_constant:
OS << "__constant";
break;
+ case LangAS::opencl_private:
+ OS << "__private";
+ break;
case LangAS::opencl_generic:
OS << "__generic";
break;
@@ -1689,6 +1693,7 @@
OS << "__attribute__((address_space(";
OS << addrspace - LangAS::FirstTargetAddressSpace;
OS << ")))";
+ }
}
}
if (Qualifiers::GC gc = getObjCGCAttr()) {
Index: lib/AST/ItaniumMangle.cpp
===================================================================
--- lib/AST/ItaniumMangle.cpp
+++ lib/AST/ItaniumMangle.cpp
@@ -2155,23 +2155,26 @@
if (Context.getASTContext().addressSpaceMapManglingFor(AS)) {
// <target-addrspace> ::= "AS" <address-space-number>
unsigned TargetAS = Context.getASTContext().getTargetAddressSpace(AS);
- ASString = "AS" + llvm::utostr(TargetAS);
+ if (TargetAS != 0)
+ ASString = "AS" + llvm::utostr(TargetAS);
} else {
switch (AS) {
default: llvm_unreachable("Not a language specific address space");
- // <OpenCL-addrspace> ::= "CL" [ "global" | "local" | "constant |
- // "generic" ]
+ // <OpenCL-addrspace> ::= "CL" [ "global" | "local" | "constant" |
+ // "private"| "generic" ]
case LangAS::opencl_global: ASString = "CLglobal"; break;
case LangAS::opencl_local: ASString = "CLlocal"; break;
case LangAS::opencl_constant: ASString = "CLconstant"; break;
+ case LangAS::opencl_private: ASString = "CLprivate"; break;
case LangAS::opencl_generic: ASString = "CLgeneric"; break;
// <CUDA-addrspace> ::= "CU" [ "device" | "constant" | "shared" ]
case LangAS::cuda_device: ASString = "CUdevice"; break;
case LangAS::cuda_constant: ASString = "CUconstant"; break;
case LangAS::cuda_shared: ASString = "CUshared"; break;
}
}
- mangleVendorQualifier(ASString);
+ if (!ASString.empty())
+ mangleVendorQualifier(ASString);
}
// The ARC ownership qualifiers start with underscores.
Index: lib/AST/Expr.cpp
===================================================================
--- lib/AST/Expr.cpp
+++ lib/AST/Expr.cpp
@@ -3250,20 +3250,20 @@
// Check that it is a cast to void*.
if (const PointerType *PT = CE->getType()->getAs<PointerType>()) {
QualType Pointee = PT->getPointeeType();
- Qualifiers Q = Pointee.getQualifiers();
- // In OpenCL v2.0 generic address space acts as a placeholder
- // and should be ignored.
- bool IsASValid = true;
- if (Ctx.getLangOpts().OpenCLVersion >= 200) {
- if (Pointee.getAddressSpace() == LangAS::opencl_generic)
- Q.removeAddressSpace();
- else
- IsASValid = false;
- }
+ // Only (void*)0 or equivalent are treated as nullptr. If pointee type
+ // has non-default address space it is not treated as nullptr.
+ bool PointeeHasDefaultAS = false;
+ if (Ctx.getLangOpts().OpenCL)
+ PointeeHasDefaultAS =
+ (Ctx.getLangOpts().OpenCLVersion >= 200 &&
+ Pointee.getAddressSpace() == LangAS::opencl_generic) ||
+ (Ctx.getLangOpts().OpenCLVersion < 200 &&
+ Pointee.getAddressSpace() == LangAS::opencl_private);
+ else
+ PointeeHasDefaultAS = Pointee.getAddressSpace() == LangAS::Default;
- if (IsASValid && !Q.hasQualifiers() &&
- Pointee->isVoidType() && // to void*
- CE->getSubExpr()->getType()->isIntegerType()) // from int.
+ if (PointeeHasDefaultAS && Pointee->isVoidType() && // to void*
+ CE->getSubExpr()->getType()->isIntegerType()) // from int.
return CE->getSubExpr()->isNullPointerConstant(Ctx, NPC);
}
}
Index: lib/AST/ASTContext.cpp
===================================================================
--- lib/AST/ASTContext.cpp
+++ lib/AST/ASTContext.cpp
@@ -707,6 +707,7 @@
1, // opencl_global
3, // opencl_local
2, // opencl_constant
+ 0, // opencl_private
4, // opencl_generic
5, // cuda_device
6, // cuda_constant
@@ -2276,10 +2277,11 @@
return QualType(eq, fastQuals);
}
-QualType
-ASTContext::getAddrSpaceQualType(QualType T, unsigned AddressSpace) const {
+QualType ASTContext::getAddrSpaceQualType(QualType T, unsigned AddressSpace,
+ bool ImplicitFlag) const {
QualType CanT = getCanonicalType(T);
- if (CanT.getAddressSpace() == AddressSpace)
+ if (CanT.getAddressSpace() == AddressSpace &&
+ CanT.getQualifiers().getImplicitAddressSpaceFlag() == ImplicitFlag)
return T;
// If we are composing extended qualifiers together, merge together
@@ -2292,6 +2294,7 @@
assert(!Quals.hasAddressSpace() &&
"Type cannot be in multiple addr spaces!");
Quals.addAddressSpace(AddressSpace);
+ Quals.setImplicitAddressSpaceFlag(ImplicitFlag);
return getExtQualType(TypeNode, Quals);
}
@@ -8082,6 +8085,7 @@
// If the qualifiers are different, the types aren't compatible... mostly.
Qualifiers LQuals = LHSCan.getLocalQualifiers();
Qualifiers RQuals = RHSCan.getLocalQualifiers();
+ RQuals.setImplicitAddressSpaceFlag(LQuals.getImplicitAddressSpaceFlag());
if (LQuals != RQuals) {
// If any of these qualifiers are different, we have a type
// mismatch.
Index: include/clang/Basic/AddressSpaces.h
===================================================================
--- include/clang/Basic/AddressSpaces.h
+++ include/clang/Basic/AddressSpaces.h
@@ -35,6 +35,7 @@
opencl_global,
opencl_local,
opencl_constant,
+ opencl_private,
opencl_generic,
// CUDA specific address spaces.
Index: include/clang/AST/Type.h
===================================================================
--- include/clang/AST/Type.h
+++ include/clang/AST/Type.h
@@ -152,8 +152,8 @@
enum {
/// The maximum supported address space number.
- /// 23 bits should be enough for anyone.
- MaxAddressSpace = 0x7fffffu,
+ /// 22 bits should be enough for anyone.
+ MaxAddressSpace = 0x3fffffu,
/// The width of the "fast" qualifier mask.
FastWidth = 3,
@@ -329,6 +329,13 @@
return (lifetime == OCL_Strong || lifetime == OCL_Weak);
}
+ bool getImplicitAddressSpaceFlag() const { return Mask & IMask; }
+ void setImplicitAddressSpaceFlag(bool Value) {
+ Mask = (Mask & ~IMask) | (((uint32_t)Value) << IShift);
+ }
+ void removeImplicitAddressSpaceFlag() {
+ setImplicitAddressSpaceFlag(false);
+ }
bool hasAddressSpace() const { return Mask & AddressSpaceMask; }
unsigned getAddressSpace() const { return Mask >> AddressSpaceShift; }
bool hasTargetSpecificAddressSpace() const {
@@ -353,7 +360,10 @@
Mask = (Mask & ~AddressSpaceMask)
| (((uint32_t) space) << AddressSpaceShift);
}
- void removeAddressSpace() { setAddressSpace(0); }
+ void removeAddressSpace() {
+ setAddressSpace(0);
+ removeImplicitAddressSpaceFlag();
+ }
void addAddressSpace(unsigned space) {
assert(space);
setAddressSpace(space);
@@ -536,20 +546,21 @@
}
private:
-
- // bits: |0 1 2|3|4 .. 5|6 .. 8|9 ... 31|
- // |C R V|U|GCAttr|Lifetime|AddressSpace|
+ // bits: |0 1 2|3|4 .. 5|6 .. 8|9|10 ... 31|
+ // |C R V|U|GCAttr|Lifetime|I|AddressSpace |
uint32_t Mask;
static const uint32_t UMask = 0x8;
static const uint32_t UShift = 3;
static const uint32_t GCAttrMask = 0x30;
static const uint32_t GCAttrShift = 4;
static const uint32_t LifetimeMask = 0x1C0;
static const uint32_t LifetimeShift = 6;
+ static const uint32_t IMask = 0x200;
+ static const uint32_t IShift = 9;
static const uint32_t AddressSpaceMask =
- ~(CVRMask | UMask | GCAttrMask | LifetimeMask);
- static const uint32_t AddressSpaceShift = 9;
+ ~(CVRMask | UMask | GCAttrMask | LifetimeMask | IMask);
+ static const uint32_t AddressSpaceShift = 10;
};
/// A std::pair-like structure for storing a qualified type split
Index: include/clang/AST/ASTContext.h
===================================================================
--- include/clang/AST/ASTContext.h
+++ include/clang/AST/ASTContext.h
@@ -1067,7 +1067,8 @@
/// The resulting type has a union of the qualifiers from T and the address
/// space. If T already has an address space specifier, it is silently
/// replaced.
- QualType getAddrSpaceQualType(QualType T, unsigned AddressSpace) const;
+ QualType getAddrSpaceQualType(QualType T, unsigned AddressSpace,
+ bool IsImplicit = false) const;
/// \brief Apply Objective-C protocol qualifiers to the given type.
/// \param allowOnPointerType specifies if we can apply protocol
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
