1// RUN: %clang_cc1 %s -verify=expected -pedantic -fsyntax-only 2// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify=expected -pedantic -fsyntax-only 3// RUN: %clang_cc1 %s -cl-std=CL3.0 -cl-ext=+__opencl_c_generic_address_space -verify=expected -pedantic -fsyntax-only 4// RUN: %clang_cc1 %s -cl-std=clc++1.0 -verify -pedantic -fsyntax-only 5// RUN: %clang_cc1 %s -cl-std=clc++2021 -cl-ext=+__opencl_c_generic_address_space -verify -pedantic -fsyntax-only 6 7__constant int ci = 1; 8 9// __constant ints are allowed in constant expressions. 10enum use_ci_in_enum { enumerator = ci }; 11typedef int use_ci_in_array_bound[ci]; 12 13// general constant folding of array bounds is not permitted 14typedef 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}} 15 16__kernel void foo(__global int *gip) { 17 __local int li; 18 __local int lj = 2; // expected-error {{'__local' variable cannot have an initializer}} 19 20 int *ip; 21#if ((!__OPENCL_CPP_VERSION__) && (__OPENCL_C_VERSION__ < 200)) 22 ip = gip; // expected-error {{assigning '__global int *__private' to '__private int *__private' changes address space of pointer}} 23 ip = &li; // expected-error {{assigning '__local int *' to '__private int *__private' changes address space of pointer}} 24 ip = &ci; // expected-error {{assigning '__constant int *' to '__private int *__private' changes address space of pointer}} 25#else 26 ip = gip; 27 ip = &li; 28 ip = &ci; 29#if !__OPENCL_CPP_VERSION__ 30// expected-error@-2 {{assigning '__constant int *' to '__generic int *__private' changes address space of pointer}} 31#else 32// expected-error@-4 {{assigning '__constant int *' to '__generic int *' changes address space of pointer}} 33#endif 34#endif 35} 36 37void explicit_cast(__global int *g, __local int *l, __constant int *c, __private int *p, const __constant int *cc) { 38 g = (__global int *)l; 39#if !__OPENCL_CPP_VERSION__ 40// expected-error@-2 {{casting '__local int *' to type '__global int *' changes address space of pointer}} 41#else 42// expected-error@-4 {{C-style cast from '__local int *' to '__global int *' converts between mismatching address spaces}} 43#endif 44 g = (__global int *)c; 45#if !__OPENCL_CPP_VERSION__ 46// expected-error@-2 {{casting '__constant int *' to type '__global int *' changes address space of pointer}} 47#else 48// expected-error@-4 {{C-style cast from '__constant int *' to '__global int *' converts between mismatching address spaces}} 49#endif 50 g = (__global int *)cc; 51#if !__OPENCL_CPP_VERSION__ 52// expected-error@-2 {{casting 'const __constant int *' to type '__global int *' changes address space of pointer}} 53#else 54// expected-error@-4 {{C-style cast from 'const __constant int *' to '__global int *' converts between mismatching address spaces}} 55#endif 56 g = (__global int *)p; 57#if !__OPENCL_CPP_VERSION__ 58// expected-error@-2 {{casting '__private int *' to type '__global int *' changes address space of pointer}} 59#else 60// expected-error@-4 {{C-style cast from '__private int *' to '__global int *' converts between mismatching address spaces}} 61#endif 62 l = (__local int *)g; 63#if !__OPENCL_CPP_VERSION__ 64// expected-error@-2 {{casting '__global int *' to type '__local int *' changes address space of pointer}} 65#else 66// expected-error@-4 {{C-style cast from '__global int *' to '__local int *' converts between mismatching address spaces}} 67#endif 68 l = (__local int *)c; 69#if !__OPENCL_CPP_VERSION__ 70// expected-error@-2 {{casting '__constant int *' to type '__local int *' changes address space of pointer}} 71#else 72// expected-error@-4 {{C-style cast from '__constant int *' to '__local int *' converts between mismatching address spaces}} 73#endif 74 l = (__local int *)cc; 75#if !__OPENCL_CPP_VERSION__ 76// expected-error@-2 {{casting 'const __constant int *' to type '__local int *' changes address space of pointer}} 77#else 78// expected-error@-4 {{C-style cast from 'const __constant int *' to '__local int *' converts between mismatching address spaces}} 79#endif 80 l = (__local int *)p; 81#if !__OPENCL_CPP_VERSION__ 82// expected-error@-2 {{casting '__private int *' to type '__local int *' changes address space of pointer}} 83#else 84// expected-error@-4 {{C-style cast from '__private int *' to '__local int *' converts between mismatching address spaces}} 85#endif 86 c = (__constant int *)g; 87#if !__OPENCL_CPP_VERSION__ 88// expected-error@-2 {{casting '__global int *' to type '__constant int *' changes address space of pointer}} 89#else 90// expected-error@-4 {{C-style cast from '__global int *' to '__constant int *' converts between mismatching address spaces}} 91#endif 92 c = (__constant int *)l; 93#if !__OPENCL_CPP_VERSION__ 94// expected-error@-2 {{casting '__local int *' to type '__constant int *' changes address space of pointer}} 95#else 96// expected-error@-4 {{C-style cast from '__local int *' to '__constant int *' converts between mismatching address spaces}} 97#endif 98 c = (__constant int *)p; 99#if !__OPENCL_CPP_VERSION__ 100// expected-error@-2 {{casting '__private int *' to type '__constant int *' changes address space of pointer}} 101#else 102// expected-error@-4 {{C-style cast from '__private int *' to '__constant int *' converts between mismatching address spaces}} 103#endif 104 p = (__private int *)g; 105#if !__OPENCL_CPP_VERSION__ 106// expected-error@-2 {{casting '__global int *' to type '__private int *' changes address space of pointer}} 107#else 108// expected-error@-4 {{C-style cast from '__global int *' to '__private int *' converts between mismatching address spaces}} 109#endif 110 p = (__private int *)l; 111#if !__OPENCL_CPP_VERSION__ 112// expected-error@-2 {{casting '__local int *' to type '__private int *' changes address space of pointer}} 113#else 114// expected-error@-4 {{C-style cast from '__local int *' to '__private int *' converts between mismatching address spaces}} 115#endif 116 p = (__private int *)c; 117#if !__OPENCL_CPP_VERSION__ 118// expected-error@-2 {{casting '__constant int *' to type '__private int *' changes address space of pointer}} 119#else 120// expected-error@-4 {{C-style cast from '__constant int *' to '__private int *' converts between mismatching address spaces}} 121#endif 122 p = (__private int *)cc; 123#if !__OPENCL_CPP_VERSION__ 124// expected-error@-2 {{casting 'const __constant int *' to type '__private int *' changes address space of pointer}} 125#else 126// expected-error@-4 {{C-style cast from 'const __constant int *' to '__private int *' converts between mismatching address spaces}} 127#endif 128} 129 130void ok_explicit_casts(__global int *g, __global int *g2, __local int *l, __local int *l2, __private int *p, __private int *p2) { 131 g = (__global int *)g2; 132 l = (__local int *)l2; 133 p = (__private int *)p2; 134} 135 136#if !__OPENCL_CPP_VERSION__ 137void nested(__global int *g, __global int * __private *gg, __local int *l, __local int * __private *ll, __global float * __private *gg_f) { 138 g = gg; // expected-error {{assigning '__global int *__private *__private' to '__global int *__private' changes address space of pointer}} 139 g = l; // expected-error {{assigning '__local int *__private' to '__global int *__private' changes address space of pointer}} 140 g = ll; // expected-error {{assigning '__local int *__private *__private' to '__global int *__private' changes address space of pointer}} 141 g = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__global int *__private' changes address space of pointer}} 142 g = (__global int *)gg_f; // expected-error {{casting '__global float *__private *' to type '__global int *' changes address space of pointer}} 143 144 gg = g; // expected-error {{assigning '__global int *__private' to '__global int *__private *__private' changes address space of pointer}} 145 gg = l; // expected-error {{assigning '__local int *__private' to '__global int *__private *__private' changes address space of pointer}} 146 gg = ll; // expected-error {{assigning '__local int *__private *__private' to '__global int *__private *__private' changes address space of nested pointer}} 147 gg = gg_f; // expected-warning {{incompatible pointer types assigning to '__global int *__private *__private' from '__global float *__private *__private'}} 148 gg = (__global int * __private *)gg_f; 149 150 l = g; // expected-error {{assigning '__global int *__private' to '__local int *__private' changes address space of pointer}} 151 l = gg; // expected-error {{assigning '__global int *__private *__private' to '__local int *__private' changes address space of pointer}} 152 l = ll; // expected-error {{assigning '__local int *__private *__private' to '__local int *__private' changes address space of pointer}} 153 l = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__local int *__private' changes address space of pointer}} 154 l = (__local int *)gg_f; // expected-error {{casting '__global float *__private *' to type '__local int *' changes address space of pointer}} 155 156 ll = g; // expected-error {{assigning '__global int *__private' to '__local int *__private *__private' changes address space of pointer}} 157 ll = gg; // expected-error {{assigning '__global int *__private *__private' to '__local int *__private *__private' changes address space of nested pointer}} 158 ll = l; // expected-error {{assigning '__local int *__private' to '__local int *__private *__private' changes address space of pointer}} 159 ll = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__local int *__private *__private' changes address space of nested pointer}} 160 ll = (__local int * __private *)gg_f; // expected-warning {{casting '__global float *__private *' to type '__local int *__private *' discards qualifiers in nested pointer types}} 161 162 gg_f = g; // expected-error {{assigning '__global int *__private' to '__global float *__private *__private' changes address space of pointer}} 163 gg_f = gg; // expected-warning {{incompatible pointer types assigning to '__global float *__private *__private' from '__global int *__private *__private'}} 164 gg_f = l; // expected-error {{assigning '__local int *__private' to '__global float *__private *__private' changes address space of pointer}} 165 gg_f = ll; // expected-error {{assigning '__local int *__private *__private' to '__global float *__private *__private' changes address space of nested pointer}} 166 gg_f = (__global float * __private *)gg; 167 168 // FIXME: This doesn't seem right. This should be an error, not a warning. 169 __local int * __global * __private * lll; 170 lll = gg; // expected-warning {{incompatible pointer types assigning to '__local int *__global *__private *__private' from '__global int *__private *__private'}} 171 172 typedef __local int * l_t; 173 typedef __global int * g_t; 174 __private l_t * pl; 175 __private g_t * pg; 176 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}} 177 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}} 178 gg = pg; 179 pg = gg; 180 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}} 181 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}} 182 183 ll = (__local int * __private *)(void *)gg; 184 void *vp = ll; 185} 186#else 187void nested(__global int *g, __global int * __private *gg, __local int *l, __local int * __private *ll, __global float * __private *gg_f) { 188 g = gg; // expected-error {{assigning '__global int *__private *__private' to '__global int *' changes address space of pointer}} 189 g = l; // expected-error {{assigning '__local int *__private' to '__global int *' changes address space of pointer}} 190 g = ll; // expected-error {{assigning '__local int *__private *__private' to '__global int *' changes address space of pointer}} 191 g = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__global int *' changes address space of pointer}} 192 g = (__global int *)gg_f; // expected-error {{C-style cast from '__global float *__private *' to '__global int *' converts between mismatching address spaces}} 193 194 gg = g; // expected-error {{assigning '__global int *__private' to '__global int *__private *' changes address space of pointer}} 195 gg = l; // expected-error {{assigning '__local int *__private' to '__global int *__private *' changes address space of pointer}} 196 gg = ll; // expected-error {{assigning '__local int *__private *__private' to '__global int *__private *' changes address space of nested pointer}} 197 gg = gg_f; // expected-error {{incompatible pointer types assigning to '__global int *__private *' from '__global float *__private *__private'}} 198 gg = (__global int * __private *)gg_f; 199 200 l = g; // expected-error {{assigning '__global int *__private' to '__local int *' changes address space of pointer}} 201 l = gg; // expected-error {{assigning '__global int *__private *__private' to '__local int *' changes address space of pointer}} 202 l = ll; // expected-error {{assigning '__local int *__private *__private' to '__local int *' changes address space of pointer}} 203 l = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__local int *' changes address space of pointer}} 204 l = (__local int *)gg_f; // expected-error {{C-style cast from '__global float *__private *' to '__local int *' converts between mismatching address spaces}} 205 206 ll = g; // expected-error {{assigning '__global int *__private' to '__local int *__private *' changes address space of pointer}} 207 ll = gg; // expected-error {{assigning '__global int *__private *__private' to '__local int *__private *' changes address space of nested pointer}} 208 ll = l; // expected-error {{assigning '__local int *__private' to '__local int *__private *' changes address space of pointer}} 209 ll = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__local int *__private *' changes address space of nested pointer}} 210 ll = (__local int *__private *)gg; //expected-warning{{C-style cast from '__global int *__private *' to '__local int *__private *' changes address space of nested pointers}} 211 212 gg_f = g; // expected-error {{assigning '__global int *__private' to '__global float *__private *' changes address space of pointer}} 213 gg_f = gg; // expected-error {{incompatible pointer types assigning to '__global float *__private *' from '__global int *__private *__private'}} 214 gg_f = l; // expected-error {{assigning '__local int *__private' to '__global float *__private *' changes address space of pointer}} 215 gg_f = ll; // expected-error {{assigning '__local int *__private *__private' to '__global float *__private *' changes address space of nested pointer}} 216 gg_f = (__global float * __private *)gg; 217 218 typedef __local int * l_t; 219 typedef __global int * g_t; 220 __private l_t * pl; 221 __private g_t * pg; 222 gg = pl; // expected-error {{assigning '__private l_t *__private' (aka '__local int *__private *__private') to '__global int *__private *' changes address space of nested pointer}} 223 pl = gg; // expected-error {{assigning '__global int *__private *__private' to '__private l_t *' (aka '__local int *__private *') changes address space of nested pointer}} 224 gg = pg; 225 pg = gg; 226 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}} 227 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}} 228 229 ll = (__local int * __private *)(void *)gg; 230 void *vp = ll; 231} 232#endif 233 234__private int func_return_priv(void); //expected-error {{return type cannot be qualified with address space}} 235__global int func_return_global(void); //expected-error {{return type cannot be qualified with address space}} 236__local int func_return_local(void); //expected-error {{return type cannot be qualified with address space}} 237__constant int func_return_constant(void); //expected-error {{return type cannot be qualified with address space}} 238#if __OPENCL_C_VERSION__ >= 200 239__generic int func_return_generic(void); //expected-error {{return type cannot be qualified with address space}} 240#endif 241 242void func_multiple_addr(void) { 243 typedef __private int private_int_t; 244 __private __local int var1; // expected-error {{multiple address spaces specified for type}} 245 __private __local int *var2; // expected-error {{multiple address spaces specified for type}} 246 __local private_int_t var3; // expected-error {{multiple address spaces specified for type}} 247 __local private_int_t *var4; // expected-error {{multiple address spaces specified for type}} 248 __private private_int_t var5; // expected-warning {{multiple identical address spaces specified for type}} 249 __private private_int_t *var6;// expected-warning {{multiple identical address spaces specified for type}} 250} 251 252void func_with_array_param(const unsigned data[16]); 253 254__kernel void k() { 255 unsigned data[16]; 256 func_with_array_param(data); 257} 258 259void func_multiple_addr2(void) { 260 typedef __private int private_int_t; 261 __attribute__((opencl_global)) __private int var1; // expected-error {{multiple address spaces specified for type}} \ 262 // expected-error {{function scope variable cannot be declared in global address space}} 263 __private __attribute__((opencl_global)) int *var2; // expected-error {{multiple address spaces specified for type}} 264 __attribute__((opencl_global)) private_int_t var3; // expected-error {{multiple address spaces specified for type}} 265 __attribute__((opencl_global)) private_int_t *var4; // expected-error {{multiple address spaces specified for type}} 266 __attribute__((opencl_private)) private_int_t var5; // expected-warning {{multiple identical address spaces specified for type}} 267 __attribute__((opencl_private)) private_int_t *var6; // expected-warning {{multiple identical address spaces specified for type}} 268#if __OPENCL_CPP_VERSION__ 269 __global int [[clang::opencl_private]] var7; // expected-error {{multiple address spaces specified for type}} 270 __global int [[clang::opencl_private]] *var8; // expected-error {{multiple address spaces specified for type}} 271 private_int_t [[clang::opencl_private]] var9; // expected-warning {{multiple identical address spaces specified for type}} 272 private_int_t [[clang::opencl_private]] *var10; // expected-warning {{multiple identical address spaces specified for type}} 273#endif // !__OPENCL_CPP_VERSION__ 274} 275