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 144 145 146 147 148 149 150 151 152 153 154 155 156 157 158 159 160 161 162 163 164 165 166 167 168 169 170 171 172 173 174 175 176 177 178 179 180 181 182 183 184 185 186 187 188 189 190 191 192 193 194 195 196 197 198 199 200 201 202 203 204 205 206 207 208 209 210 211 212 213 214 215 216 217 218 219 220 221 222 223 224 225 226 227 228 229 230 231 232 233 234 235 236 237 238 239 240 241 242 243 244 245 246 247 248 249 250 251 252 253 254 255 256 257 258 259 260 261 262 263 264 265 266 267 268 269 270 271 272 273
|
// 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=CL3.0 -cl-ext=+__opencl_c_generic_address_space -verify -pedantic -fsyntax-only
// RUN: %clang_cc1 %s -cl-std=clc++ -verify -pedantic -fsyntax-only
__constant int ci = 1;
// __constant ints are allowed in constant expressions.
enum use_ci_in_enum { enumerator = ci };
typedef int use_ci_in_array_bound[ci];
// general constant folding of array bounds is not permitted
typedef int folding_in_array_bounds[&ci + 3 - &ci]; // expected-error-re {{{{variable length arrays are not supported in OpenCL|array size is not a constant expression}}}} expected-note {{cannot refer to element 3}}
__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 *__private' to '__private int *__private' changes address space of pointer}}
ip = &li; // expected-error {{assigning '__local int *' to '__private int *__private' changes address space of pointer}}
ip = &ci; // expected-error {{assigning '__constant int *' to '__private int *__private' 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 *__private' changes address space of pointer}}
#else
// expected-error@-4 {{assigning '__constant int *' to '__generic int *' changes address space of pointer}}
#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 '__private int *' to type '__global int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from '__private 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 '__private int *' to type '__local int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from '__private 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 '__private int *' to type '__constant int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from '__private 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 '__private int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from '__global int *' to '__private int *' converts between mismatching address spaces}}
#endif
p = (__private int *)l;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{casting '__local int *' to type '__private int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from '__local int *' to '__private int *' converts between mismatching address spaces}}
#endif
p = (__private int *)c;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{casting '__constant int *' to type '__private int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from '__constant int *' to '__private int *' converts between mismatching address spaces}}
#endif
p = (__private int *)cc;
#if !__OPENCL_CPP_VERSION__
// expected-error@-2 {{casting 'const __constant int *' to type '__private int *' changes address space of pointer}}
#else
// expected-error@-4 {{C-style cast from 'const __constant int *' to '__private 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;
}
#if !__OPENCL_CPP_VERSION__
void nested(__global int *g, __global int * __private *gg, __local int *l, __local int * __private *ll, __global float * __private *gg_f) {
g = gg; // expected-error {{assigning '__global int *__private *__private' to '__global int *__private' changes address space of pointer}}
g = l; // expected-error {{assigning '__local int *__private' to '__global int *__private' changes address space of pointer}}
g = ll; // expected-error {{assigning '__local int *__private *__private' to '__global int *__private' changes address space of pointer}}
g = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__global int *__private' changes address space of pointer}}
g = (__global int *)gg_f; // expected-error {{casting '__global float *__private *' to type '__global int *' changes address space of pointer}}
gg = g; // expected-error {{assigning '__global int *__private' to '__global int *__private *__private' changes address space of pointer}}
gg = l; // expected-error {{assigning '__local int *__private' to '__global int *__private *__private' changes address space of pointer}}
gg = ll; // expected-error {{assigning '__local int *__private *__private' to '__global int *__private *__private' changes address space of nested pointer}}
gg = gg_f; // expected-warning {{incompatible pointer types assigning to '__global int *__private *__private' from '__global float *__private *__private'}}
gg = (__global int * __private *)gg_f;
l = g; // expected-error {{assigning '__global int *__private' to '__local int *__private' changes address space of pointer}}
l = gg; // expected-error {{assigning '__global int *__private *__private' to '__local int *__private' changes address space of pointer}}
l = ll; // expected-error {{assigning '__local int *__private *__private' to '__local int *__private' changes address space of pointer}}
l = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__local int *__private' changes address space of pointer}}
l = (__local int *)gg_f; // expected-error {{casting '__global float *__private *' to type '__local int *' changes address space of pointer}}
ll = g; // expected-error {{assigning '__global int *__private' to '__local int *__private *__private' changes address space of pointer}}
ll = gg; // expected-error {{assigning '__global int *__private *__private' to '__local int *__private *__private' changes address space of nested pointer}}
ll = l; // expected-error {{assigning '__local int *__private' to '__local int *__private *__private' changes address space of pointer}}
ll = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__local int *__private *__private' changes address space of nested pointer}}
ll = (__local int * __private *)gg_f; // expected-warning {{casting '__global float *__private *' to type '__local int *__private *' discards qualifiers in nested pointer types}}
gg_f = g; // expected-error {{assigning '__global int *__private' to '__global float *__private *__private' changes address space of pointer}}
gg_f = gg; // expected-warning {{incompatible pointer types assigning to '__global float *__private *__private' from '__global int *__private *__private'}}
gg_f = l; // expected-error {{assigning '__local int *__private' to '__global float *__private *__private' changes address space of pointer}}
gg_f = ll; // expected-error {{assigning '__local int *__private *__private' to '__global float *__private *__private' changes address space of nested pointer}}
gg_f = (__global float * __private *)gg;
// FIXME: This doesn't seem right. This should be an error, not a warning.
__local int * __global * __private * lll;
lll = gg; // expected-warning {{incompatible pointer types assigning to '__local int *__global *__private *__private' from '__global int *__private *__private'}}
typedef __local int * l_t;
typedef __global int * g_t;
__private l_t * pl;
__private g_t * pg;
gg = pl; // expected-error {{assigning '__private l_t *__private' (aka '__local int *__private *__private') to '__global int *__private *__private' changes address space of nested pointer}}
pl = gg; // expected-error {{assigning '__global int *__private *__private' to '__private l_t *__private' (aka '__local int *__private *__private') changes address space of nested pointer}}
gg = pg;
pg = gg;
pg = pl; // expected-error {{assigning '__private l_t *__private' (aka '__local int *__private *__private') to '__private g_t *__private' (aka '__global int *__private *__private') changes address space of nested pointer}}
pl = pg; // expected-error {{assigning '__private g_t *__private' (aka '__global int *__private *__private') to '__private l_t *__private' (aka '__local int *__private *__private') changes address space of nested pointer}}
ll = (__local int * __private *)(void *)gg;
void *vp = ll;
}
#else
void nested(__global int *g, __global int * __private *gg, __local int *l, __local int * __private *ll, __global float * __private *gg_f) {
g = gg; // expected-error {{assigning '__global int *__private *__private' to '__global int *' changes address space of pointer}}
g = l; // expected-error {{assigning '__local int *__private' to '__global int *' changes address space of pointer}}
g = ll; // expected-error {{assigning '__local int *__private *__private' to '__global int *' changes address space of pointer}}
g = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__global int *' changes address space of pointer}}
g = (__global int *)gg_f; // expected-error {{C-style cast from '__global float *__private *' to '__global int *' converts between mismatching address spaces}}
gg = g; // expected-error {{assigning '__global int *__private' to '__global int *__private *' changes address space of pointer}}
gg = l; // expected-error {{assigning '__local int *__private' to '__global int *__private *' changes address space of pointer}}
gg = ll; // expected-error {{assigning '__local int *__private *__private' to '__global int *__private *' changes address space of nested pointer}}
gg = gg_f; // expected-error {{incompatible pointer types assigning to '__global int *__private *' from '__global float *__private *__private'}}
gg = (__global int * __private *)gg_f;
l = g; // expected-error {{assigning '__global int *__private' to '__local int *' changes address space of pointer}}
l = gg; // expected-error {{assigning '__global int *__private *__private' to '__local int *' changes address space of pointer}}
l = ll; // expected-error {{assigning '__local int *__private *__private' to '__local int *' changes address space of pointer}}
l = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__local int *' changes address space of pointer}}
l = (__local int *)gg_f; // expected-error {{C-style cast from '__global float *__private *' to '__local int *' converts between mismatching address spaces}}
ll = g; // expected-error {{assigning '__global int *__private' to '__local int *__private *' changes address space of pointer}}
ll = gg; // expected-error {{assigning '__global int *__private *__private' to '__local int *__private *' changes address space of nested pointer}}
ll = l; // expected-error {{assigning '__local int *__private' to '__local int *__private *' changes address space of pointer}}
ll = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__local int *__private *' changes address space of nested pointer}}
ll = (__local int *__private *)gg; //expected-warning{{C-style cast from '__global int *__private *' to '__local int *__private *' changes address space of nested pointers}}
gg_f = g; // expected-error {{assigning '__global int *__private' to '__global float *__private *' changes address space of pointer}}
gg_f = gg; // expected-error {{incompatible pointer types assigning to '__global float *__private *' from '__global int *__private *__private'}}
gg_f = l; // expected-error {{assigning '__local int *__private' to '__global float *__private *' changes address space of pointer}}
gg_f = ll; // expected-error {{assigning '__local int *__private *__private' to '__global float *__private *' changes address space of nested pointer}}
gg_f = (__global float * __private *)gg;
typedef __local int * l_t;
typedef __global int * g_t;
__private l_t * pl;
__private g_t * pg;
gg = pl; // expected-error {{assigning '__private l_t *__private' (aka '__local int *__private *__private') to '__global int *__private *' changes address space of nested pointer}}
pl = gg; // expected-error {{assigning '__global int *__private *__private' to '__private l_t *' (aka '__local int *__private *') changes address space of nested pointer}}
gg = pg;
pg = gg;
pg = pl; // expected-error {{assigning '__private l_t *__private' (aka '__local int *__private *__private') to '__private g_t *' (aka '__global int *__private *') changes address space of nested pointer}}
pl = pg; // expected-error {{assigning '__private g_t *__private' (aka '__global int *__private *__private') to '__private l_t *' (aka '__local int *__private *') changes address space of nested pointer}}
ll = (__local int * __private *)(void *)gg;
void *vp = ll;
}
#endif
__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}}
}
void func_with_array_param(const unsigned data[16]);
__kernel void k() {
unsigned data[16];
func_with_array_param(data);
}
void func_multiple_addr2(void) {
typedef __private int private_int_t;
__private __attribute__((opencl_global)) int var1; // expected-error {{multiple address spaces specified for type}} \
// expected-error {{function scope variable cannot be declared in global address space}}
__private __attribute__((opencl_global)) int *var2; // expected-error {{multiple address spaces specified for type}}
__attribute__((opencl_global)) private_int_t var3; // expected-error {{multiple address spaces specified for type}}
__attribute__((opencl_global)) private_int_t *var4; // expected-error {{multiple address spaces specified for type}}
__attribute__((opencl_private)) private_int_t var5; // expected-warning {{multiple identical address spaces specified for type}}
__attribute__((opencl_private)) private_int_t *var6; // expected-warning {{multiple identical address spaces specified for type}}
#if __OPENCL_CPP_VERSION__
[[clang::opencl_private]] __global int var7; // expected-error {{multiple address spaces specified for type}}
[[clang::opencl_private]] __global int *var8; // expected-error {{multiple address spaces specified for type}}
[[clang::opencl_private]] private_int_t var9; // expected-warning {{multiple identical address spaces specified for type}}
[[clang::opencl_private]] private_int_t *var10; // expected-warning {{multiple identical address spaces specified for type}}
#endif // !__OPENCL_CPP_VERSION__
}
|