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 | |