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