blob: f5fa43d9222810c3221b81fd46188f08bb248e0f (
plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
|
// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -pedantic -fsyntax-only
// RUN: %clang_cc1 %s -cl-std=c++ -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_CPP_VERSION__) && (__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;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{assigning '__constant int *' to '__generic int *' changes address space of pointer}}
#else
// expected-error@-4 {{assigning to '__generic int *' from incompatible type '__constant int *'}}
#endif
#endif
}
void explicit_cast(__global int *g, __local int *l, __constant int *c, __private int *p, const __constant int *cc) {
g = (__global int *)l;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{casting '__local int *' to type '__global int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from '__local int *' to '__global int *' converts between mismatching address spaces}}
#endif
g = (__global int *)c;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{casting '__constant int *' to type '__global int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from '__constant int *' to '__global int *' converts between mismatching address spaces}}
#endif
g = (__global int *)cc;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{casting 'const __constant int *' to type '__global int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from 'const __constant int *' to '__global int *' converts between mismatching address spaces}}
#endif
g = (__global int *)p;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{casting 'int *' to type '__global int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from 'int *' to '__global int *' converts between mismatching address spaces}}
#endif
l = (__local int *)g;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{casting '__global int *' to type '__local int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from '__global int *' to '__local int *' converts between mismatching address spaces}}
#endif
l = (__local int *)c;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{casting '__constant int *' to type '__local int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from '__constant int *' to '__local int *' converts between mismatching address spaces}}
#endif
l = (__local int *)cc;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{casting 'const __constant int *' to type '__local int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from 'const __constant int *' to '__local int *' converts between mismatching address spaces}}
#endif
l = (__local int *)p;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{casting 'int *' to type '__local int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from 'int *' to '__local int *' converts between mismatching address spaces}}
#endif
c = (__constant int *)g;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{casting '__global int *' to type '__constant int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from '__global int *' to '__constant int *' converts between mismatching address spaces}}
#endif
c = (__constant int *)l;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{casting '__local int *' to type '__constant int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from '__local int *' to '__constant int *' converts between mismatching address spaces}}
#endif
c = (__constant int *)p;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{casting 'int *' to type '__constant int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from 'int *' to '__constant int *' converts between mismatching address spaces}}
#endif
p = (__private int *)g;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{casting '__global int *' to type 'int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from '__global int *' to 'int *' converts between mismatching address spaces}}
#endif
p = (__private int *)l;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{casting '__local int *' to type 'int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from '__local int *' to 'int *' converts between mismatching address spaces}}
#endif
p = (__private int *)c;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{casting '__constant int *' to type 'int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from '__constant int *' to 'int *' converts between mismatching address spaces}}
#endif
p = (__private int *)cc;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{casting 'const __constant int *' to type 'int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from 'const __constant int *' to 'int *' converts between mismatching address spaces}}
#endif
}
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;
__private __local int var1; // expected-error {{multiple address spaces specified for type}}
__private __local 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}}
__private private_int_t var5; // expected-warning {{multiple identical address spaces specified for type}}
__private private_int_t *var6;// expected-warning {{multiple identical address spaces specified for type}}
}
|
