clang/test/SemaOpenCL/address-spaces.cl

   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.
  10 enum use_ci_in_enum { enumerator = ci };
  11 typedef int use_ci_in_array_bound[ci];
  12
  13 // general constant folding of array bounds is not permitted
  14 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}}
  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
  37 void 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
 130 void 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__
 137 void 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
 187 void 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
 242 void 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
 252 void 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
 259 void 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 }