275 lines · plain
1// RUN: %clang_cc1 %s -verify=expected -pedantic -fsyntax-only2// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify=expected -pedantic -fsyntax-only3// RUN: %clang_cc1 %s -cl-std=CL3.0 -cl-ext=+__opencl_c_generic_address_space -verify=expected -pedantic -fsyntax-only4// RUN: %clang_cc1 %s -cl-std=clc++1.0 -verify -pedantic -fsyntax-only5// RUN: %clang_cc1 %s -cl-std=clc++2021 -cl-ext=+__opencl_c_generic_address_space -verify -pedantic -fsyntax-only6 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 permitted14typedef 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#else26 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#else32// expected-error@-4 {{assigning '__constant int *' to '__generic int *' changes address space of pointer}}33#endif34#endif35}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#else42// expected-error@-4 {{C-style cast from '__local int *' to '__global int *' converts between mismatching address spaces}}43#endif44 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#else48// expected-error@-4 {{C-style cast from '__constant int *' to '__global int *' converts between mismatching address spaces}}49#endif50 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#else54// expected-error@-4 {{C-style cast from 'const __constant int *' to '__global int *' converts between mismatching address spaces}}55#endif56 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#else60// expected-error@-4 {{C-style cast from '__private int *' to '__global int *' converts between mismatching address spaces}}61#endif62 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#else66// expected-error@-4 {{C-style cast from '__global int *' to '__local int *' converts between mismatching address spaces}}67#endif68 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#else72// expected-error@-4 {{C-style cast from '__constant int *' to '__local int *' converts between mismatching address spaces}}73#endif74 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#else78// expected-error@-4 {{C-style cast from 'const __constant int *' to '__local int *' converts between mismatching address spaces}}79#endif80 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#else84// expected-error@-4 {{C-style cast from '__private int *' to '__local int *' converts between mismatching address spaces}}85#endif86 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#else90// expected-error@-4 {{C-style cast from '__global int *' to '__constant int *' converts between mismatching address spaces}}91#endif92 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#else96// expected-error@-4 {{C-style cast from '__local int *' to '__constant int *' converts between mismatching address spaces}}97#endif98 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#else102// expected-error@-4 {{C-style cast from '__private int *' to '__constant int *' converts between mismatching address spaces}}103#endif104 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#else108// expected-error@-4 {{C-style cast from '__global int *' to '__private int *' converts between mismatching address spaces}}109#endif110 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#else114// expected-error@-4 {{C-style cast from '__local int *' to '__private int *' converts between mismatching address spaces}}115#endif116 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#else120// expected-error@-4 {{C-style cast from '__constant int *' to '__private int *' converts between mismatching address spaces}}121#endif122 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#else126// expected-error@-4 {{C-style cast from 'const __constant int *' to '__private int *' converts between mismatching address spaces}}127#endif128}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-error {{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-error {{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 __local int * __global * __private * lll;169 lll = gg; // expected-error {{incompatible pointer types assigning to '__local int *__global *__private *__private' from '__global int *__private *__private'}}170 171 typedef __local int * l_t;172 typedef __global int * g_t;173 __private l_t * pl;174 __private g_t * pg;175 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}}176 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}}177 gg = pg;178 pg = gg;179 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}}180 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}}181 182 ll = (__local int * __private *)(void *)gg;183 void *vp = ll;184}185#else186void nested(__global int *g, __global int * __private *gg, __local int *l, __local int * __private *ll, __global float * __private *gg_f) {187 g = gg; // expected-error {{assigning '__global int *__private *__private' to '__global int *' changes address space of pointer}}188 g = l; // expected-error {{assigning '__local int *__private' to '__global int *' changes address space of pointer}}189 g = ll; // expected-error {{assigning '__local int *__private *__private' to '__global int *' changes address space of pointer}}190 g = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__global int *' changes address space of pointer}}191 g = (__global int *)gg_f; // expected-error {{C-style cast from '__global float *__private *' to '__global int *' converts between mismatching address spaces}}192 193 gg = g; // expected-error {{assigning '__global int *__private' to '__global int *__private *' changes address space of pointer}}194 gg = l; // expected-error {{assigning '__local int *__private' to '__global int *__private *' changes address space of pointer}}195 gg = ll; // expected-error {{assigning '__local int *__private *__private' to '__global int *__private *' changes address space of nested pointer}}196 gg = gg_f; // expected-error {{incompatible pointer types assigning to '__global int *__private *' from '__global float *__private *__private'}}197 gg = (__global int * __private *)gg_f;198 199 l = g; // expected-error {{assigning '__global int *__private' to '__local int *' changes address space of pointer}}200 l = gg; // expected-error {{assigning '__global int *__private *__private' to '__local int *' changes address space of pointer}}201 l = ll; // expected-error {{assigning '__local int *__private *__private' to '__local int *' changes address space of pointer}}202 l = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__local int *' changes address space of pointer}}203 l = (__local int *)gg_f; // expected-error {{C-style cast from '__global float *__private *' to '__local int *' converts between mismatching address spaces}}204 205 ll = g; // expected-error {{assigning '__global int *__private' to '__local int *__private *' changes address space of pointer}}206 ll = gg; // expected-error {{assigning '__global int *__private *__private' to '__local int *__private *' changes address space of nested pointer}}207 ll = l; // expected-error {{assigning '__local int *__private' to '__local int *__private *' changes address space of pointer}}208 ll = gg_f; // expected-error {{assigning '__global float *__private *__private' to '__local int *__private *' changes address space of nested pointer}}209 ll = (__local int *__private *)gg; //expected-warning{{C-style cast from '__global int *__private *' to '__local int *__private *' changes address space of nested pointers}}210 211 gg_f = g; // expected-error {{assigning '__global int *__private' to '__global float *__private *' changes address space of pointer}}212 gg_f = gg; // expected-error {{incompatible pointer types assigning to '__global float *__private *' from '__global int *__private *__private'}}213 gg_f = l; // expected-error {{assigning '__local int *__private' to '__global float *__private *' changes address space of pointer}}214 gg_f = ll; // expected-error {{assigning '__local int *__private *__private' to '__global float *__private *' changes address space of nested pointer}}215 gg_f = (__global float * __private *)gg;216 217 typedef __local int * l_t;218 typedef __global int * g_t;219 __private l_t * pl;220 __private g_t * pg;221 gg = pl; // expected-error {{assigning '__private l_t *__private' (aka '__local int *__private *__private') to '__global int *__private *' changes address space of nested pointer}}222 pl = gg; // expected-error {{assigning '__global int *__private *__private' to '__private l_t *' (aka '__local int *__private *') changes address space of nested pointer}}223 gg = pg;224 pg = gg;225 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}}226 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}}227 228 ll = (__local int * __private *)(void *)gg;229 void *vp = ll;230}231#endif232 233__private int func_return_priv(void); //expected-error {{return type cannot be qualified with address space}}234__global int func_return_global(void); //expected-error {{return type cannot be qualified with address space}}235__local int func_return_local(void); //expected-error {{return type cannot be qualified with address space}}236__constant int func_return_constant(void); //expected-error {{return type cannot be qualified with address space}}237#if __OPENCL_C_VERSION__ >= 200238__generic int func_return_generic(void); //expected-error {{return type cannot be qualified with address space}}239#endif240 241void func_multiple_addr(void) {242 typedef __private int private_int_t;243 __private __local int var1; // expected-error {{multiple address spaces specified for type}}244 __private __local int *var2; // expected-error {{multiple address spaces specified for type}}245 __local private_int_t var3; // expected-error {{multiple address spaces specified for type}}246 __local private_int_t *var4; // expected-error {{multiple address spaces specified for type}}247 __private private_int_t var5; // expected-warning {{multiple identical address spaces specified for type}}248 __private private_int_t *var6;// expected-warning {{multiple identical address spaces specified for type}}249}250 251void func_with_array_param(const unsigned data[16]);252 253__kernel void k() {254 unsigned data[16];255 func_with_array_param(data);256}257 258void func_multiple_addr2(void) {259 typedef __private int private_int_t;260 __attribute__((opencl_global)) __private int var1; // expected-error {{multiple address spaces specified for type}} \261 // expected-error {{function scope variable cannot be declared in global address space}}262 __private __attribute__((opencl_global)) int *var2; // expected-error {{multiple address spaces specified for type}}263 __attribute__((opencl_global)) private_int_t var3; // expected-error {{multiple address spaces specified for type}}264 __attribute__((opencl_global)) private_int_t *var4; // expected-error {{multiple address spaces specified for type}}265 __attribute__((opencl_private)) private_int_t var5; // expected-warning {{multiple identical address spaces specified for type}}266 __attribute__((opencl_private)) private_int_t *var6; // expected-warning {{multiple identical address spaces specified for type}}267#if __OPENCL_CPP_VERSION__268 __global int [[clang::opencl_private]] var7; // expected-error {{multiple address spaces specified for type}} \269 // expected-error {{function scope variable cannot be declared in global address space}}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