yaxunl updated this revision to Diff 117020.
yaxunl marked 3 inline comments as done.
yaxunl added a comment.
Revised by Anastasia's comments.
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/CodeGen/blocks-opencl.cl
test/CodeGenOpenCL/address-spaces-mangling.cl
test/CodeGenOpenCL/address-spaces.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/extern.cl
test/SemaOpenCL/invalid-block.cl
test/SemaOpenCL/invalid-pipes-cl2.0.cl
test/SemaOpenCL/null_literal.cl
test/SemaOpenCL/storageclass-cl20.cl
test/SemaOpenCL/storageclass.cl
Index: test/SemaOpenCL/storageclass.cl
===================================================================
--- test/SemaOpenCL/storageclass.cl
+++ test/SemaOpenCL/storageclass.cl
@@ -5,6 +5,20 @@
int G3 = 0; // expected-error{{program scope variable must reside in constant address space}}
global int G4 = 0; // expected-error{{program scope variable must reside in constant address space}}
+static float g_implicit_static_var = 0; // expected-error {{program scope variable must reside in constant address space}}
+static constant float g_constant_static_var = 0;
+static global float g_global_static_var = 0; // expected-error {{program scope variable must reside in constant address space}}
+static local float g_local_static_var = 0; // expected-error {{program scope variable must reside in constant address space}}
+static private float g_private_static_var = 0; // expected-error {{program scope variable must reside in constant address space}}
+static generic float g_generic_static_var = 0; // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{program scope variable must reside in constant address space}}
+
+extern float g_implicit_extern_var; // expected-error {{extern variable must reside in constant address space}}
+extern constant float g_constant_extern_var;
+extern global float g_global_extern_var; // expected-error {{extern variable must reside in constant address space}}
+extern local float g_local_extern_var; // expected-error {{extern variable must reside in constant address space}}
+extern private float g_private_extern_var; // expected-error {{extern variable must reside in constant address space}}
+extern generic float g_generic_extern_var; // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{extern variable must reside in constant address space}}
+
void kernel foo(int x) {
// static is not allowed at local scope before CL2.0
static int S1 = 5; // expected-error{{variables in function scope cannot be declared static}}
@@ -45,10 +59,17 @@
__attribute__((address_space(100))) int L4; // expected-error{{automatic variable qualified with an invalid address space}}
}
+ static float l_implicit_static_var = 0; // expected-error {{variables in function scope cannot be declared static}}
+ static constant float l_constant_static_var = 0; // expected-error {{variables in function scope cannot be declared static}}
+ static global float l_global_static_var = 0; // expected-error {{variables in function scope cannot be declared static}}
+ static local float l_local_static_var = 0; // expected-error {{variables in function scope cannot be declared static}}
+ static private float l_private_static_var = 0; // expected-error {{variables in function scope cannot be declared static}}
+ static generic float l_generic_static_var = 0; // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{variables in function scope cannot be declared static}}
- extern constant float L5;
- extern local float L6; // expected-error{{extern variable must reside in constant address space}}
-
- static int L7 = 0; // expected-error{{variables in function scope cannot be declared static}}
- static int L8; // expected-error{{variables in function scope cannot be declared static}}
+ extern float l_implicit_extern_var; // expected-error {{extern variable must reside in constant address space}}
+ extern constant float l_constant_extern_var;
+ extern global float l_global_extern_var; // expected-error {{extern variable must reside in constant address space}}
+ extern local float l_local_extern_var; // expected-error {{extern variable must reside in constant address space}}
+ extern private float l_private_extern_var; // expected-error {{extern variable must reside in constant address space}}
+ extern generic float l_generic_extern_var; // expected-error{{OpenCL version 1.2 does not support the 'generic' type qualifier}} // expected-error {{extern variable must reside in constant address space}}
}
Index: test/SemaOpenCL/storageclass-cl20.cl
===================================================================
--- test/SemaOpenCL/storageclass-cl20.cl
+++ test/SemaOpenCL/storageclass-cl20.cl
@@ -1,21 +1,41 @@
// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL2.0
-static constant int G1 = 0;
int G2 = 0;
global int G3 = 0;
local int G4 = 0; // expected-error{{program scope variable must reside in global or constant address space}}
-void kernel foo() {
- static int S1 = 5;
- static global int S2 = 5;
- static private int S3 = 5; // expected-error{{static local variable must reside in global or constant address space}}
+static float g_implicit_static_var = 0;
+static constant float g_constant_static_var = 0;
+static global float g_global_static_var = 0;
+static local float g_local_static_var = 0; // expected-error {{program scope variable must reside in global or constant address space}}
+static private float g_private_static_var = 0; // expected-error {{program scope variable must reside in global or constant address space}}
+static generic float g_generic_static_var = 0; // expected-error {{program scope variable must reside in global or constant address space}}
+
+extern float g_implicit_extern_var;
+extern constant float g_constant_extern_var;
+extern global float g_global_extern_var;
+extern local float g_local_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
+extern private float g_private_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
+extern generic float g_generic_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
+void kernel foo() {
constant int L1 = 0;
local int L2;
global int L3; // expected-error{{function scope variable cannot be declared in global address space}}
generic int L4; // expected-error{{automatic variable qualified with an invalid address space}}
__attribute__((address_space(100))) int L5; // expected-error{{automatic variable qualified with an invalid address space}}
- extern global int G5;
- extern int G6; // expected-error{{extern variable must reside in global or constant address space}}
+ static float l_implicit_static_var = 0;
+ static constant float l_constant_static_var = 0;
+ static global float l_global_static_var = 0;
+ static local float l_local_static_var = 0; // expected-error {{static local variable must reside in global or constant address space}}
+ static private float l_private_static_var = 0; // expected-error {{static local variable must reside in global or constant address space}}
+ static generic float l_generic_static_var = 0; // expected-error {{static local variable must reside in global or constant address space}}
+
+ extern float l_implicit_extern_var;
+ extern constant float l_constant_extern_var;
+ extern global float l_global_extern_var;
+ extern local float l_local_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
+ extern private float l_private_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
+ extern generic float l_generic_extern_var; // expected-error {{extern variable must reside in global or constant address space}}
}
Index: test/SemaOpenCL/null_literal.cl
===================================================================
--- test/SemaOpenCL/null_literal.cl
+++ test/SemaOpenCL/null_literal.cl
@@ -1,29 +1,68 @@
// 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)
void foo(){
+ global int *g1 = NULL;
+ global int *g2 = (global void *)0;
+ global int *g3 = (constant void *)0; // expected-error{{initializing '__global int *' with an expression of type '__constant void *' changes address space of pointer}}
+ global int *g4 = (local void *)0; // expected-error{{initializing '__global int *' with an expression of type '__local void *' changes address space of pointer}}
+ global int *g5 = (private void *)0; // expected-error{{initializing '__global int *' with an expression of type '__private void *' changes address space of pointer}}
-global int* ptr1 = NULL;
+ constant int *c1 = NULL;
+ constant int *c2 = (global void *)0; // expected-error{{initializing '__constant int *' with an expression of type '__global void *' changes address space of pointer}}
+ constant int *c3 = (constant void *)0;
+ constant int *c4 = (local void *)0; // expected-error{{initializing '__constant int *' with an expression of type '__local void *' changes address space of pointer}}
+ constant int *c5 = (private void *)0; // expected-error{{initializing '__constant int *' with an expression of type '__private void *' changes address space of pointer}}
-global int* ptr2 = (global void*)0;
+ local int *l1 = NULL;
+ local int *l2 = (global void *)0; // expected-error{{initializing '__local int *' with an expression of type '__global void *' changes address space of pointer}}
+ local int *l3 = (constant void *)0; // expected-error{{initializing '__local int *' with an expression of type '__constant void *' changes address space of pointer}}
+ local int *l4 = (local void *)0;
+ local int *l5 = (private void *)0; // expected-error{{initializing '__local int *' with an expression of type '__private void *' changes address space of pointer}}
-constant int* ptr3 = NULL;
+ private int *p1 = NULL;
+ private int *p2 = (global void *)0; // expected-error{{initializing '__private int *' with an expression of type '__global void *' changes address space of pointer}}
+ private int *p3 = (constant void *)0; // expected-error{{initializing '__private int *' with an expression of type '__constant void *' changes address space of pointer}}
+ private int *p4 = (local void *)0; // expected-error{{initializing '__private int *' with an expression of type '__local void *' changes address space of pointer}}
+ private int *p5 = (private void *)0;
-constant int* ptr4 = (global void*)0; // expected-error{{initializing '__constant int *' with an expression of type '__global void *' changes address space of pointer}}
+#if __OPENCL_C_VERSION__ >= 200
+ // Assigning a pointer to a pointer to narrower address space causes an error unless there is an valid explicit cast.
+ global int *g6 = (generic void *)0; // expected-error{{initializing '__global int *' with an expression of type '__generic void *' changes address space of pointer}}
+ constant int *c6 = (generic void *)0; // expected-error{{initializing '__constant int *' with an expression of type '__generic void *' changes address space of pointer}}
+ local int *l6 = (generic void *)0; // expected-error{{initializing '__local int *' with an expression of type '__generic void *' changes address space of pointer}}
+ private int *p6 = (generic void *)0; // expected-error{{initializing '__private int *' with an expression of type '__generic void *' changes address space of pointer}}
-#ifdef CL20
-// Accept explicitly pointer to generic address space in OpenCL v2.0.
-global int* ptr5 = (generic 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}}
+ global int *g7 = (global void*)(generic void *)0;
+ constant int *c7 = (constant void*)(generic void *)0; //expected-error{{casting '__generic void *' to type '__constant void *' changes address space of pointer}}
+ local int *l7 = (local void*)(generic void *)0;
+ private int *p7 = (private void*)(generic void *)0;
-bool cmp = ptr1 == NULL;
+ generic int *ge1 = NULL;
+ generic int *ge2 = (global void *)0;
+ generic int *ge3 = (constant void *)0; // expected-error{{initializing '__generic int *' with an expression of type '__constant void *' changes address space of pointer}}
+ generic int *ge4 = (local void *)0;
+ generic int *ge5 = (private void *)0;
+ generic int *ge6 = (generic void *)0;
+#endif
-cmp = ptr1 == (local void*)0; // expected-error{{comparison between ('__global int *' and '__local void *') which are pointers to non-overlapping address spaces}}
+ bool cmp;
+ cmp = g1 == NULL;
+ cmp = g1 == (global void *)0;
+ cmp = g1 == (constant void *)0; // expected-error{{comparison between ('__global int *' and '__constant void *') which are pointers to non-overlapping address spaces}}
+ cmp = g1 == (local void *)0; // expected-error{{comparison between ('__global int *' and '__local void *') which are pointers to non-overlapping address spaces}}
+ cmp = g1 == (private void *)0; // expected-error{{comparison between ('__global int *' and '__private void *') which are pointers to non-overlapping address spaces}}
-cmp = ptr3 == NULL;
+#if __OPENCL_C_VERSION__ >= 200
+ cmp = g1 == (generic void *)0;
+ cmp = ge1 == NULL;
+ cmp = ge1 == (global void *)0;
+ cmp = ge1 == (constant void *)0; // expected-error{{comparison between ('__generic int *' and '__constant void *') which are pointers to non-overlapping address spaces}}
+ cmp = ge1 == (local void *)0;
+ cmp = ge1 == (private void *)0;
+ cmp = ge1 == (generic void *)0;
+#endif
}
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/extern.cl
===================================================================
--- test/SemaOpenCL/extern.cl
+++ /dev/null
@@ -1,9 +0,0 @@
-// RUN: %clang_cc1 -x cl -cl-opt-disable -cl-std=CL1.2 -emit-llvm -ffake-address-space-map %s -o - -verify | FileCheck %s
-// expected-no-diagnostics
-
-// CHECK: @foo = external addrspace(2) constant float
-extern constant float foo;
-
-kernel void test(global float* buf) {
- buf[0] += foo;
-}
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.cl
===================================================================
--- test/CodeGenOpenCL/address-spaces.cl
+++ test/CodeGenOpenCL/address-spaces.cl
@@ -7,6 +7,24 @@
// RUN: %clang_cc1 %s -O0 -triple amdgcn-mesa-mesa3d -emit-llvm -o - | FileCheck --check-prefixes=CHECK,SPIR %s
// RUN: %clang_cc1 %s -O0 -triple r600-- -emit-llvm -o - | FileCheck --check-prefixes=CHECK,SPIR %s
+// SPIR: %struct.S = type { i32, i32, i32* }
+// CL20SPIR: %struct.S = type { i32, i32, i32 addrspace(4)* }
+struct S {
+ int x;
+ int y;
+ int *z;
+};
+
+// CL20-DAG: @g_extern_var = external addrspace(1) global float
+// CL20-DAG: @l_extern_var = external addrspace(1) global float
+// CL20-DAG: @test_static.l_static_var = internal addrspace(1) global float 0.000000e+00
+// CL20-DAG: @g_static_var = internal addrspace(1) global float 0.000000e+00
+
+#ifdef CL20
+// CL20-DAG: @g_s = common addrspace(1) global %struct.S zeroinitializer
+struct S g_s;
+#endif
+
// SPIR: i32* %arg
// GIZ: i32 addrspace(5)* %arg
void f__p(__private int *arg) {}
@@ -58,3 +76,52 @@
// CL20-DAG: @f.ii = internal addrspace(1) global i32 0
#endif
}
+
+typedef int int_td;
+typedef int *intp_td;
+// SPIR: define void @test_typedef(i32 addrspace(1)* %x, i32 addrspace(2)* %y, i32* %z)
+void test_typedef(global int_td *x, constant int_td *y, intp_td z) {
+ *x = *y;
+ *z = 0;
+}
+
+// SPIR: define void @test_struct()
+void test_struct() {
+ // SPIR: %ps = alloca %struct.S*
+ // CL20SPIR: %ps = alloca %struct.S addrspace(4)*
+ struct S *ps;
+ // SPIR: store i32 0, i32* %x
+ // CL20SPIR: store i32 0, i32 addrspace(4)* %x
+ ps->x = 0;
+#ifdef CL20
+ // CL20SPIR: store i32 0, i32 addrspace(1)* getelementptr inbounds (%struct.S, %struct.S addrspace(1)* @g_s, i32 0, i32 0)
+ g_s.x = 0;
+#endif
+}
+
+// SPIR-LABEL: define void @test_void_par()
+void test_void_par(void) {}
+
+// SPIR-LABEL: define i32 @test_func_return_type()
+int test_func_return_type(void) {
+ return 0;
+}
+
+#ifdef CL20
+extern float g_extern_var;
+
+// CL20-LABEL: define {{.*}}void @test_extern(
+kernel void test_extern(global float *buf) {
+ extern float l_extern_var;
+ buf[0] += g_extern_var + l_extern_var;
+}
+
+static float g_static_var;
+
+// CL20-LABEL: define {{.*}}void @test_static(
+kernel void test_static(global float *buf) {
+ static float l_static_var;
+ buf[0] += g_static_var + l_static_var;
+}
+
+#endif
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: test/CodeGen/blocks-opencl.cl
===================================================================
--- test/CodeGen/blocks-opencl.cl
+++ test/CodeGen/blocks-opencl.cl
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -O0 %s -ffake-address-space-map -emit-llvm -o - -fblocks -triple x86_64-unknown-unknown | FileCheck %s
+// RUN: %clang_cc1 -cl-std=CL2.0 -O0 %s -ffake-address-space-map -emit-llvm -o - -fblocks -triple x86_64-unknown-unknown | FileCheck %s
// This used to crash due to trying to generate a bitcase from a cstring
// in the constant address space to i8* in AS0.
Index: lib/Sema/SemaType.cpp
===================================================================
--- lib/Sema/SemaType.cpp
+++ lib/Sema/SemaType.cpp
@@ -4834,7 +4834,6 @@
TypeSourceInfo *ReturnTypeInfo = nullptr;
QualType T = GetDeclSpecTypeForDeclarator(state, ReturnTypeInfo);
-
if (D.isPrototypeContext() && getLangOpts().ObjCAutoRefCount)
inferARCWriteback(state, T);
@@ -5574,7 +5573,7 @@
ASIdx = LangAS::opencl_generic; break;
default:
assert(Attr.getKind() == AttributeList::AT_OpenCLPrivateAddressSpace);
- ASIdx = 0; break;
+ ASIdx = LangAS::opencl_private; break;
}
}
@@ -6806,6 +6805,92 @@
}
}
+static void deduceOpenCLImplicitAddrSpace(TypeProcessingState &State,
+ QualType &T, TypeAttrLocation TAL) {
+ Declarator &D = State.getDeclarator();
+
+ // Handle the cases where address space should not be deduced.
+ //
+ // The pointee type of a pointer type is alwasy deduced since a pointer always
+ // points to some memory location which should has an address space.
+ //
+ // There are situations that at the point of certain declarations, the address
+ // space may be unknown and better to be left as default. For example, when
+ // definining a typedef or struct type, they are not associated with any
+ // specific address space. Later on, they may be used with any address space
+ // to declare a variable.
+ //
+ // The return value of a function is r-value, therefore should not have
+ // address space.
+ //
+ // The void type does not occupy memory, therefore should not have address
+ // space, except when it is used as a pointee type.
+ //
+ // Since LLVM assumes function type is in default address space, it should not
+ // have address space.
+ auto ChunkIndex = State.getCurrentChunkIndex();
+ bool IsPointee =
+ ChunkIndex > 0 &&
+ (D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Pointer ||
+ D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::BlockPointer);
+ bool IsFuncReturnType =
+ ChunkIndex > 0 &&
+ D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Function;
+ bool IsFuncType =
+ ChunkIndex < D.getNumTypeObjects() &&
+ D.getTypeObject(ChunkIndex).Kind == DeclaratorChunk::Function;
+ if ( // Do not deduce addr space for function return type and function type,
+ // otherwise it will fail some sema check.
+ IsFuncReturnType || IsFuncType ||
+ // Do not deduce addr space for member types of struct, except the pointee
+ // type of a pointer member type.
+ (D.getContext() == Declarator::MemberContext && !IsPointee) ||
+ // Do not deduce addr space for types used to define a typedef and the
+ // typedef itself, except the pointee type of a pointer type which is used
+ // to define the typedef.
+ (D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_typedef &&
+ !IsPointee) ||
+ // Do not deduce addr space of the void type, e.g. in f(void), otherwise
+ // it will fail some sema check.
+ (T->isVoidType() && !IsPointee))
+ return;
+
+ unsigned ImpAddr;
+ // Put OpenCL automatic variable in private address space.
+ // OpenCL v1.2 s6.5:
+ // The default 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.
+ if (State.getSema().getLangOpts().OpenCLVersion <= 120) {
+ ImpAddr = LangAS::opencl_private;
+ } else {
+ // If address space is not set, OpenCL 2.0 defines non private default
+ // address spaces for some cases:
+ // OpenCL 2.0, section 6.5:
+ // The address space for a variable at program scope or a static variable
+ // inside a function can either be __global or __constant, but defaults to
+ // __global if not specified.
+ // (...)
+ // Pointers that are declared without pointing to a named address space
+ // point to the generic address space.
+ if (IsPointee) {
+ ImpAddr = LangAS::opencl_generic;
+ } else {
+ if (D.getContext() == Declarator::FileContext) {
+ ImpAddr = LangAS::opencl_global;
+ } else {
+ if (D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_static ||
+ D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_extern) {
+ ImpAddr = LangAS::opencl_global;
+ } else {
+ ImpAddr = LangAS::opencl_private;
+ }
+ }
+ }
+ }
+ T = State.getSema().Context.getAddrSpaceQualType(T, ImpAddr, true);
+}
+
static void processTypeAttrs(TypeProcessingState &state, QualType &type,
TypeAttrLocation TAL, AttributeList *attrs) {
// Scan through and apply attributes to this type where it makes sense. Some
@@ -6977,39 +7062,11 @@
}
}
- // If address space is not set, OpenCL 2.0 defines non private default
- // address spaces for some cases:
- // OpenCL 2.0, section 6.5:
- // The address space for a variable at program scope or a static variable
- // inside a function can either be __global or __constant, but defaults to
- // __global if not specified.
- // (...)
- // 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 &&
- (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);
- } 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 = state.getSema().Context.getAddrSpaceQualType(
- type, LangAS::opencl_global);
- }
+ if (!state.getSema().getLangOpts().OpenCL ||
+ type.getAddressSpace() != LangAS::Default)
+ return;
+
+ deduceOpenCLImplicitAddrSpace(state, type, TAL);
}
void Sema::completeExprArrayBound(Expr *E) {
Index: lib/Sema/SemaDecl.cpp
===================================================================
--- lib/Sema/SemaDecl.cpp
+++ lib/Sema/SemaDecl.cpp
@@ -6246,7 +6246,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();
}
@@ -7349,7 +7349,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();
@@ -7984,7 +7984,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;
}
@@ -8754,9 +8755,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();
@@ -11866,13 +11865,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
@@ -340,7 +340,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'";
@@ -784,8 +784,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
@@ -1662,11 +1662,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;
@@ -1677,6 +1678,9 @@
case LangAS::cuda_constant:
OS << "__constant";
break;
+ case LangAS::opencl_private:
+ OS << "__private";
+ break;
case LangAS::opencl_generic:
OS << "__generic";
break;
@@ -1691,6 +1695,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
@@ -2219,23 +2219,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
@@ -3282,20 +3282,16 @@
// 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;
- }
-
- if (IsASValid && !Q.hasQualifiers() &&
- Pointee->isVoidType() && // to void*
- CE->getSubExpr()->getType()->isIntegerType()) // from int.
+ // Only (void*)0 or equivalent are treated as nullptr. If pointee type
+ // has non-default address space it is not treated as nullptr.
+ // (__generic void*)0 in OpenCL 2.0 should not be treated as nullptr
+ // since it cannot be assigned to a pointer to constant address space.
+ bool PointeeHasDefaultAS =
+ Pointee.getAddressSpace() == LangAS::Default ||
+ Pointee.getQualifiers().getImplicitAddressSpaceFlag();
+
+ 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
@@ -2282,10 +2283,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
@@ -2298,6 +2300,7 @@
assert(!Quals.hasAddressSpace() &&
"Type cannot be in multiple addr spaces!");
Quals.addAddressSpace(AddressSpace);
+ Quals.setImplicitAddressSpaceFlag(ImplicitFlag);
return getExtQualType(TypeNode, Quals);
}
@@ -8102,6 +8105,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,17 @@
return (lifetime == OCL_Strong || lifetime == OCL_Weak);
}
+ /// True if the non-default address space is not explicit in the source
+ /// code but deduced by context. This flag is used when printing
+ /// types or performing semantic checks if the explicity of an address
+ /// space makes difference.
+ 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 +364,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 +550,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
@@ -1068,7 +1068,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
[email protected]
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
