1 /* Test OpenACC 'kernels' construct decomposition. */
2
3 /* { dg-additional-options "--param=openacc-kernels=decompose" } */
4
5 /* { dg-additional-options "-fopt-info-all-omp" }
6 { dg-additional-options "-foffload=-fopt-info-all-omp" } */
7
8 /* { dg-additional-options "--param=openacc-privatization=noisy" }
9 { dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
10 Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
11 { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
12
13 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
14 passed to 'incr' may be unset, and in that case, it will be set to [...]",
15 so to maintain compatibility with earlier Tcl releases, we manually
16 initialize counter variables:
17 { dg-line l_dummy[variable c_compute 0 c_loop_c 0 c_loop_i 0] }
18 { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
19 "WARNING: dg-line var l_dummy defined, but not used". */
20
21 #undef NDEBUG
22 #include <assert.h>
23
24 static int g1;
25 static int g2;
26
27 /* PR100280, etc. */
28
29 static void f1 ()
30 {
31 int a = 0;
32 #define N 123
33 int b[N] = { 0 };
34 unsigned long long f1;
35
36 #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
37 /* { dg-note {OpenACC 'kernels' decomposition: variable 'f1' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
38 { dg-note {variable 'f1' made addressable} {} { target *-*-* } l_compute$c_compute } */
39 /* { dg-note {OpenACC 'kernels' decomposition: variable 'a' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
40 { dg-note {variable 'a' made addressable} {} { target *-*-* } l_compute$c_compute } */
41 /* { dg-note {OpenACC 'kernels' decomposition: variable 'g2' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
42 { dg-note {variable 'g2' made addressable} {} { target *-*-* } l_compute$c_compute } */
43 /* { dg-note {OpenACC 'kernels' decomposition: variable 'g1' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
44 { dg-note {variable 'g1' made addressable} {} { target *-*-* } l_compute$c_compute } */
45 {
46 /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
47 int c = 234;
48 /* { dg-note {OpenACC 'kernels' decomposition: variable 'c' declared in block requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
49 { dg-note {variable 'c' made addressable} {} { target *-*-* } l_compute$c_compute }
50 { dg-note {variable 'c' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */
51
52 #pragma acc loop independent gang /* { dg-line l_loop_i[incr c_loop_i] } */
53 /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { target *-*-* } l_loop_i$c_loop_i } */
54 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
55 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop_i$c_loop_i } */
56 /* { dg-optimized "assigned OpenACC gang loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
57 for (int i = 0; i < N; ++i)
58 b[i] = c;
59
60 /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
61 a = c;
62
63 /* PR104132, PR104133, PR104774 */
64 {
65 /* Use the 'kernels'-top-level 'int c' as loop variable. */
66
67 #pragma acc loop /* { dg-line l_loop_c[incr c_loop_c] } */
68 /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_c$c_loop_c } */
69 /* { dg-note {variable 'c' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_c$c_loop_c } */
70 /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_c$c_loop_c } */
71 for (c = 0; c < N / 2; c++)
72 b[c] -= 10;
73
74 #pragma acc loop /* { dg-line l_loop_c[incr c_loop_c] } */
75 /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_c$c_loop_c } */
76 /* { dg-note {variable 'c' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_c$c_loop_c } */
77 /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_c$c_loop_c } */
78 for (c = 0; c < N / 2; c++)
79 g1 = c;
80
81 #pragma acc loop /* { dg-line l_loop_c[incr c_loop_c] } */
82 /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_c$c_loop_c } */
83 /* { dg-note {variable 'c' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_c$c_loop_c } */
84 /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_c$c_loop_c } */
85 for (c = 0; c <= N; c++)
86 g2 += c;
87 /* { dg-note {variable 'g2\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
88
89 /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
90 f1 = 1;
91 #pragma acc loop /* { dg-line l_loop_c[incr c_loop_c] } */
92 /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_c$c_loop_c } */
93 /* { dg-note {variable 'c' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_c$c_loop_c } */
94 /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_c$c_loop_c } */
95 for (c = 20; c > 0; --c)
96 f1 *= c;
97
98 {
99 /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
100 unsigned long long f2 = 1;
101 /* { dg-note {OpenACC 'kernels' decomposition: variable 'f2' declared in block requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
102 { dg-note {variable 'f2' made addressable} {} { target *-*-* } l_compute$c_compute }
103 { dg-note {variable 'f2' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */
104 #pragma acc loop independent reduction(*: f2) /* { dg-line l_loop_c[incr c_loop_c] } */
105 /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { target *-*-* } l_loop_c$c_loop_c } */
106 /* { dg-note {variable 'c' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_c$c_loop_c } */
107 /* { dg-optimized {assigned OpenACC gang vector loop parallelism} {} { target *-*-* } l_loop_c$c_loop_c } */
108 for (c = 20; c > 0; --c)
109 f2 *= c;
110
111 {
112 /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
113 if (f2 != f1)
114 __builtin_abort ();
115
116 /* As this is still in the preceding 'parloops' part:
117 { dg-bogus {note: beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
118 unsigned long long f3 = f2;
119 /* { dg-note {OpenACC 'kernels' decomposition: variable 'f3' declared in block requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
120 { dg-note {variable 'f3' made addressable} {} { target *-*-* } l_compute$c_compute }
121 { dg-note {variable 'f3' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */
122 #pragma acc loop seq /* { dg-line l_loop_c[incr c_loop_c] } */
123 /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { target *-*-* } l_loop_c$c_loop_c } */
124 /* { dg-note {variable 'c' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_c$c_loop_c } */
125 /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_c$c_loop_c } */
126 for (c = 20; c > 0; --c)
127 f3 /= c;
128
129 /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
130 if (f3 != 1)
131 __builtin_abort ();
132 }
133
134 /* As this is still in the preceding 'parloops' part:
135 { dg-bogus {note: beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
136 if (f2 != f1)
137 __builtin_abort ();
138 }
139
140 /* As this is still in the preceding 'parloops' part:
141 { dg-bogus {note: beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
142 if (c != 234)
143 __builtin_abort ();
144 /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
145 }
146 }
147
148 assert (a == 234);
149 for (int i = 0; i < N; ++i)
150 if (i < N / 2)
151 assert (b[i] == 234 - 10);
152 else
153 assert (b[i] == 234);
154 assert (g1 == N / 2 - 1);
155 assert (g2 == N * (N + 1) / 2);
156 assert (f1 == 2432902008176640000ULL);
157
158 #undef N
159 }
160
161
162 /* PR104086 */
163
164 static void f2 ()
165 {
166 #pragma acc data
167 /* { dg-bogus {note: variable [^\n\r]+ candidate for adjusting OpenACC privatization level} {TODO 'data'} { xfail *-*-* } .-1 } */
168 {
169 int i;
170
171 #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
172 /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
173 { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */
174 /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
175 i = 1;
176
177 assert (i == 1);
178
179 #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
180 /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
181 { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */
182 /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
183 i = -1;
184
185 assert (i == -1);
186 }
187
188
189 int ia[1];
190
191 #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
192 /* { dg-note {OpenACC 'kernels' decomposition: variable 'ia' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
193 { dg-note {variable 'ia' made addressable} {} { target *-*-* } l_compute$c_compute } */
194 /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
195 ia[0] = -2;
196
197 assert (ia[0] == -2);
198
199 #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
200 /* { dg-note {OpenACC 'kernels' decomposition: variable 'ia' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
201 { dg-note {variable 'ia' already made addressable} {} { target *-*-* } l_compute$c_compute } */
202 /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
203 { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute }
204 { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */
205 /* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute$c_compute } */
206 /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
207 for (int i = 0; i < 100; ++i)
208 ++ia[0];
209
210 assert (ia[0] == -2 + 100);
211 }
212
213
214 int main()
215 {
216 f1 ();
217
218 f2 ();
219
220 return 0;
221 }