1 /* Test OpenACC 'kernels' construct decomposition. */
2
3 /* { dg-additional-options "-fopt-info-omp-all" } */
4
5 /* { dg-additional-options "--param=openacc-kernels=decompose" }
6 /* { dg-additional-options "-O2" } for 'parloops'. */
7
8 /* { dg-additional-options "--param=openacc-privatization=noisy" }
9 Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
10 { dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
11
12 /* { dg-additional-options "-Wopenacc-parallelism" } for testing/documenting
13 aspects of that functionality. */
14
15 /* See also '../../gfortran.dg/goacc/kernels-decompose-2.f95'. */
16
17 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
18 passed to 'incr' may be unset, and in that case, it will be set to [...]",
19 so to maintain compatibility with earlier Tcl releases, we manually
20 initialize counter variables:
21 { dg-line l_dummy[variable c_compute 0 c_loop_i 0 c_loop_j 0 c_loop_k 0 c_part 0] }
22 { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
23 "WARNING: dg-line var l_dummy defined, but not used". */
24
25 #pragma acc routine gang
26 extern int
27 f_g (int);
28
29 #pragma acc routine worker
30 extern int
31 f_w (int);
32
33 #pragma acc routine vector
34 extern int
35 f_v (int);
36
37 #pragma acc routine seq
38 extern int
39 f_s (int);
40
41 int
42 main ()
43 {
44 int x, y, z;
45 #define N 10
46 int a[N], b[N], c[N];
47
48 #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
49 /* { dg-note {OpenACC 'kernels' decomposition: variable 'z' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
50 { dg-note {variable 'z' made addressable} {} { target *-*-* } l_compute$c_compute } */
51 /* { dg-note {OpenACC 'kernels' decomposition: variable 'y' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
52 { dg-note {variable 'y' made addressable} {} { target *-*-* } l_compute$c_compute } */
53 /* { dg-note {OpenACC 'kernels' decomposition: variable 'x' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
54 { dg-note {variable 'x' made addressable} {} { target *-*-* } l_compute$c_compute } */
55 /* { dg-note {variable 'x\.0' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
56 {
57 /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
58 x = 0;
59 y = x < 10;
60 z = x++;
61 ;
62 }
63
64 {
65 int i;
66 #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
67 /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
68 { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */
69 /* { dg-optimized {assigned OpenACC gang loop parallelism} {} { target *-*-* } l_compute$c_compute } */
70 /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
71 for (i = 0; i < N; i++)
72 a[i] = 0;
73 }
74
75 #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
76 /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } l_compute$c_compute } */
77 /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
78 { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */
79 /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */
80 {
81 int i;
82 }
83
84 #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
85 /* { dg-note {OpenACC 'kernels' decomposition: variable 'i' declared in block requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
86 { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */
87 /* { dg-note {variable 'i' declared in block is candidate for adjusting OpenACC privatization level} {} { target *-*-* } l_compute$c_compute } */
88 /* { dg-optimized {assigned OpenACC gang loop parallelism} {} { target *-*-* } l_compute$c_compute } */
89 /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
90 for (int i = 0; i < N; i++)
91 a[i] = 0;
92
93 #pragma acc kernels loop /* { dg-line l_loop_i[incr c_loop_i] } */
94 /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */
95 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
96 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
97 /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
98 for (int i = 0; i < N; i++)
99 b[i] = a[N - i - 1];
100
101 #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
102 /* { dg-note {OpenACC 'kernels' decomposition: variable 'z' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
103 { dg-note {variable 'z' already made addressable} {} { target *-*-* } l_compute$c_compute } */
104 {
105 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
106 /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */
107 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
108 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
109 /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
110 for (int i = 0; i < N; i++)
111 b[i] = a[N - i - 1];
112
113 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
114 /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */
115 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
116 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
117 /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
118 for (int i = 0; i < N; i++)
119 c[i] = a[i] * b[i];
120
121 /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
122 a[z] = 0;
123
124 #pragma acc loop /* { dg-line l_loop_i[incr c_loop_i] } */
125 /* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_loop_i$c_loop_i } */
126 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
127 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
128 /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
129 for (int i = 0; i < N; i++)
130 c[i] += a[i];
131
132 #pragma acc loop seq /* { dg-line l_loop_i[incr c_loop_i] } */
133 /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { target *-*-* } l_loop_i$c_loop_i } */
134 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
135 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
136 /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
137 for (int i = 0 + 1; i < N; i++)
138 c[i] += c[i - 1];
139 }
140
141 #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
142 /* { dg-note {OpenACC 'kernels' decomposition: variable 'y' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
143 { dg-note {variable 'y' already made addressable} {} { target *-*-* } l_compute$c_compute } */
144 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute$c_compute } */
145 /*TODO What does this mean?
146 TODO { dg-optimized "assigned OpenACC worker vector loop parallelism" "" { target *-*-* } l_compute$c_compute } */
147 {
148 #pragma acc loop independent /* { dg-line l_loop_i[incr c_loop_i] } */
149 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
150 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
151 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
152 /* { dg-optimized "assigned OpenACC gang loop parallelism" "" { target *-*-* } l_loop_i$c_loop_i } */
153 /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { target *-*-* } l_loop_i$c_loop_i } */
154 for (int i = 0; i < N; ++i)
155 #pragma acc loop independent /* { dg-line l_loop_j[incr c_loop_j] } */
156 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
157 /* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
158 /* { dg-optimized "assigned OpenACC worker loop parallelism" "" { target *-*-* } l_loop_j$c_loop_j } */
159 for (int j = 0; j < N; ++j)
160 #pragma acc loop independent /* { dg-line l_loop_k[incr c_loop_k] } */
161 /* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k$c_loop_k } */
162 /* { dg-warning "insufficient partitioning available to parallelize loop" "" { target *-*-* } l_loop_k$c_loop_k } */
163 /* { dg-optimized "assigned OpenACC seq loop parallelism" "" { target *-*-* } l_loop_k$c_loop_k } */
164 for (int k = 0; k < N; ++k)
165 a[(i + j + k) % N]
166 = b[j]
167 + f_v (c[k]); /* { dg-optimized "assigned OpenACC vector loop parallelism" } */
168
169 /*TODO Should the following turn into "gang-single" instead of "parloops"?
170 TODO The problem is that the first STMT is 'if (y <= 4) goto <D.2547>; else goto <D.2548>;', thus "parloops". */
171 /* { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
172 if (y < 5)
173 #pragma acc loop independent /* { dg-line l_loop_j[incr c_loop_j] } */
174 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
175 /* { dg-missed "unparallelized loop nest in OpenACC 'kernels' region: it's executed conditionally" "" { target *-*-* } l_loop_j$c_loop_j } */
176 for (int j = 0; j < N; ++j)
177 b[j] = f_w (c[j]);
178 }
179
180 #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
181 /* { dg-note {OpenACC 'kernels' decomposition: variable 'y' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
182 { dg-note {variable 'y' already made addressable} {} { target *-*-* } l_compute$c_compute } */
183 /* { dg-bogus "warning: region contains gang partitioned code but is not gang partitioned" "TODO 'kernels'" { xfail *-*-* } l_compute$c_compute } */
184 {
185 y = f_g (a[5]); /* { dg-line l_part[incr c_part] } */
186 /*TODO If such a construct is placed in its own part (like it is, here), can't this actually use gang paralelism, instead of "gang-single"?
187 { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } l_part$c_part } */
188 /* { dg-optimized "assigned OpenACC gang worker vector loop parallelism" "" { target *-*-* } l_part$c_part } */
189
190 #pragma acc loop independent /* { dg-line l_loop_j[incr c_loop_j] } */
191 /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { target *-*-* } l_loop_j$c_loop_j } */
192 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
193 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
194 /* { dg-optimized "assigned OpenACC gang loop parallelism" "" { target *-*-* } l_loop_j$c_loop_j } */
195 for (int j = 0; j < N; ++j)
196 b[j] = y + f_w (c[j]); /* { dg-optimized "assigned OpenACC worker vector loop parallelism" } */
197 }
198
199 #pragma acc kernels /* { dg-line l_compute[incr c_compute] } */
200 /* { dg-note {OpenACC 'kernels' decomposition: variable 'z' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
201 { dg-note {variable 'z' already made addressable} {} { target *-*-* } l_compute$c_compute } */
202 /* { dg-note {OpenACC 'kernels' decomposition: variable 'y' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute }
203 { dg-note {variable 'y' already made addressable} {} { target *-*-* } l_compute$c_compute } */
204 {
205 /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
206 y = 3;
207
208 #pragma acc loop independent /* { dg-line l_loop_j[incr c_loop_j] } */
209 /* { dg-note {parallelized loop nest in OpenACC 'kernels' region} {} { target *-*-* } l_loop_j$c_loop_j } */
210 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
211 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
212 /* { dg-optimized "assigned OpenACC gang worker loop parallelism" "" { target *-*-* } l_loop_j$c_loop_j } */
213 for (int j = 0; j < N; ++j)
214 b[j] = y + f_v (c[j]); /* { dg-optimized "assigned OpenACC vector loop parallelism" } */
215
216 /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
217 z = 2;
218 }
219
220 /* { dg-note {beginning 'gang-single' part in OpenACC 'kernels' region} {} { target *-*-* } .+1 } */
221 #pragma acc kernels
222 ;
223
224 return 0;
225 }