3 /* { dg-additional-options "-fopt-info-note-omp" }
4 { dg-additional-options "-foffload=-fopt-info-note-omp" } */
6 /* { dg-additional-options "--param=openacc-privatization=noisy" }
7 { dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
9 /* { dg-additional-options "-Wuninitialized" } */
12 { dg-xfail-run-if TODO { openacc_radeon_accel_selected && { ! __OPTIMIZE__ } } } */
14 /* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
15 passed to 'incr' may be unset, and in that case, it will be set to [...]",
16 so to maintain compatibility with earlier Tcl releases, we manually
17 initialize counter variables:
18 { dg-line l_dummy[variable c_compute 0 c_loop 0] }
19 { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
20 "WARNING: dg-line var l_dummy defined, but not used". */
29 #pragma acc parallel copyout(res) num_gangs(64) /* { dg-line l_compute[incr c_compute] } */
30 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
31 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
32 /* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
33 But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
34 No longer having address taken: tmpvar
35 Now a gimple register: tmpvar
36 However, 'tmpvar' remains in the candidate set:
37 { dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
38 Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
39 { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_compute$c_compute }
40 For nvptx offloading however, we first mark up 'tmpvar', and then later apply the privatization change -- or, with optimizations enabled, don't, because we then don't actually call 'expand_var_decl'.
41 { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_compute$c_compute }
42 { dg-bogus {note: variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_compute$c_compute }
44 /* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
49 #pragma acc loop collapse(2) gang private(tmpref) /* { dg-line l_loop[incr c_loop] } */
50 /* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
51 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
52 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
53 for (i = 0; i < 256; i++)
55 for (j = 0; j < 256; j++)
57 tmpref = (i * 256 + j) * 97;
58 res[i * 256 + j] = tmpref;
63 for (i = 0; i < 65536; i++)
73 #pragma acc parallel copyout(res) num_gangs(64) num_workers(64) /* { dg-line l_compute[incr c_compute] } */
74 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
75 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
76 /* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
77 But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
78 No longer having address taken: tmpvar
79 Now a gimple register: tmpvar
80 However, 'tmpvar' remains in the candidate set:
81 { dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
82 Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
83 { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_compute$c_compute }
84 For nvptx offloading however, we first mark up 'tmpvar', and then later apply the privatization change -- or, with optimizations enabled, don't, because we then don't actually call 'expand_var_decl'.
85 { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_compute$c_compute }
86 { dg-bogus {note: variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_compute$c_compute }
88 /* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
89 /* { dg-warning "using .num_workers \\(32\\)., ignoring 64" "" { target openacc_nvidia_accel_selected } l_compute$c_compute } */
94 #pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
95 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
96 for (i = 0; i < 256; i++)
98 #pragma acc loop worker private(tmpref) /* { dg-line l_loop[incr c_loop] } */
99 /* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
100 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
101 for (j = 0; j < 256; j++)
103 tmpref = (i * 256 + j) * 99;
104 res[i * 256 + j] = tmpref;
109 for (i = 0; i < 65536; i++)
110 if (res[i] != i * 99)
119 #pragma acc parallel copyout(res) num_gangs(64) num_workers(64) /* { dg-line l_compute[incr c_compute] } */
120 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
121 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
122 /* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
123 But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
124 No longer having address taken: tmpvar
125 Now a gimple register: tmpvar
126 However, 'tmpvar' remains in the candidate set:
127 { dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
128 Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
129 { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_compute$c_compute }
130 For nvptx offloading however, we first mark up 'tmpvar', and then later apply the privatization change -- or, with optimizations enabled, don't, because we then don't actually call 'expand_var_decl'.
131 { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_compute$c_compute }
132 { dg-bogus {note: variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_compute$c_compute }
134 /* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
135 /* { dg-warning "using .num_workers \\(32\\)., ignoring 64" "" { target openacc_nvidia_accel_selected } l_compute$c_compute } */
139 int &tmpref = tmpvar;
140 #pragma acc loop gang worker /* { dg-line l_loop[incr c_loop] } */
141 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
142 for (i = 0; i < 256; i++)
144 #pragma acc loop vector private(tmpref) /* { dg-line l_loop[incr c_loop] } */
145 /* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
146 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
147 for (j = 0; j < 256; j++)
149 tmpref = (i * 256 + j) * 101;
150 res[i * 256 + j] = tmpref;
155 for (i = 0; i < 65536; i++)
156 if (res[i] != i * 101)
160 void gangs_workers_vectors (void)
165 #pragma acc parallel copyout(res) num_gangs(64) num_workers(64) /* { dg-line l_compute[incr c_compute] } */
166 /* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
167 /* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
168 /* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
169 But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
170 No longer having address taken: tmpvar
171 Now a gimple register: tmpvar
172 However, 'tmpvar' remains in the candidate set:
173 { dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
174 Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
175 { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_compute$c_compute }
176 For nvptx offloading however, we first mark up 'tmpvar', and then later apply the privatization change -- or, with optimizations enabled, don't, because we then don't actually call 'expand_var_decl'.
177 { dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_compute$c_compute }
178 { dg-bogus {note: variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_compute$c_compute }
180 /* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
181 /* { dg-warning "using .num_workers \\(32\\)., ignoring 64" "" { target openacc_nvidia_accel_selected } l_compute$c_compute } */
185 int &tmpref = tmpvar;
186 #pragma acc loop collapse(2) gang worker vector private(tmpref) /* { dg-line l_loop[incr c_loop] } */
187 /* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
188 /* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
189 /* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
190 for (i = 0; i < 256; i++)
192 for (j = 0; j < 256; j++)
194 tmpref = (i * 256 + j) * 103;
195 res[i * 256 + j] = tmpref;
200 for (i = 0; i < 65536; i++)
201 if (res[i] != i * 103)
205 int main (int argc, char *argv[])
210 gangs_workers_vectors ();