1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
|
/* { dg-do run } */
/* { dg-additional-options "-fopt-info-note-omp" }
{ dg-additional-options "-foffload=-fopt-info-note-omp" } */
/* { dg-additional-options "--param=openacc-privatization=noisy" }
{ dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
/* { dg-additional-options "-Wuninitialized" } */
/*TODO
{ dg-xfail-run-if TODO { openacc_radeon_accel_selected && { ! __OPTIMIZE__ } } } */
/* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
passed to 'incr' may be unset, and in that case, it will be set to [...]",
so to maintain compatibility with earlier Tcl releases, we manually
initialize counter variables:
{ dg-line l_dummy[variable c_compute 0 c_loop 0] }
{ dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
"WARNING: dg-line var l_dummy defined, but not used". */
#include <stdlib.h>
void gangs (void)
{
double res[65536];
int i;
#pragma acc parallel copyout(res) num_gangs(64) /* { dg-line l_compute[incr c_compute] } */
/* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
/* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
/* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
No longer having address taken: tmpvar
Now a gimple register: tmpvar
However, 'tmpvar' remains in the candidate set:
{ dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_compute$c_compute }
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'.
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_compute$c_compute }
{ dg-bogus {note: variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_compute$c_compute }
*/
/* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
{
int i, j;
int tmpvar;
int &tmpref = tmpvar;
#pragma acc loop collapse(2) gang private(tmpref) /* { dg-line l_loop[incr c_loop] } */
/* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
/* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
for (i = 0; i < 256; i++)
{
for (j = 0; j < 256; j++)
{
tmpref = (i * 256 + j) * 97;
res[i * 256 + j] = tmpref;
}
}
}
for (i = 0; i < 65536; i++)
if (res[i] != i * 97)
abort ();
}
void workers (void)
{
double res[65536];
int i;
#pragma acc parallel copyout(res) num_gangs(64) num_workers(64) /* { dg-line l_compute[incr c_compute] } */
/* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
/* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
/* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
No longer having address taken: tmpvar
Now a gimple register: tmpvar
However, 'tmpvar' remains in the candidate set:
{ dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_compute$c_compute }
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'.
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_compute$c_compute }
{ dg-bogus {note: variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_compute$c_compute }
*/
/* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
/* { dg-warning "using .num_workers \\(32\\)., ignoring 64" "" { target openacc_nvidia_accel_selected } l_compute$c_compute } */
{
int i, j;
int tmpvar;
int &tmpref = tmpvar;
#pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
/* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
for (i = 0; i < 256; i++)
{
#pragma acc loop worker private(tmpref) /* { dg-line l_loop[incr c_loop] } */
/* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
for (j = 0; j < 256; j++)
{
tmpref = (i * 256 + j) * 99;
res[i * 256 + j] = tmpref;
}
}
}
for (i = 0; i < 65536; i++)
if (res[i] != i * 99)
abort ();
}
void vectors (void)
{
double res[65536];
int i;
#pragma acc parallel copyout(res) num_gangs(64) num_workers(64) /* { dg-line l_compute[incr c_compute] } */
/* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
/* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
/* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
No longer having address taken: tmpvar
Now a gimple register: tmpvar
However, 'tmpvar' remains in the candidate set:
{ dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_compute$c_compute }
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'.
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_compute$c_compute }
{ dg-bogus {note: variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_compute$c_compute }
*/
/* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
/* { dg-warning "using .num_workers \\(32\\)., ignoring 64" "" { target openacc_nvidia_accel_selected } l_compute$c_compute } */
{
int i, j;
int tmpvar;
int &tmpref = tmpvar;
#pragma acc loop gang worker /* { dg-line l_loop[incr c_loop] } */
/* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
for (i = 0; i < 256; i++)
{
#pragma acc loop vector private(tmpref) /* { dg-line l_loop[incr c_loop] } */
/* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
for (j = 0; j < 256; j++)
{
tmpref = (i * 256 + j) * 101;
res[i * 256 + j] = tmpref;
}
}
}
for (i = 0; i < 65536; i++)
if (res[i] != i * 101)
abort ();
}
void gangs_workers_vectors (void)
{
double res[65536];
int i;
#pragma acc parallel copyout(res) num_gangs(64) num_workers(64) /* { dg-line l_compute[incr c_compute] } */
/* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
/* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
/* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
No longer having address taken: tmpvar
Now a gimple register: tmpvar
However, 'tmpvar' remains in the candidate set:
{ dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_compute$c_compute }
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'.
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_compute$c_compute }
{ dg-bogus {note: variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_compute$c_compute }
*/
/* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
/* { dg-warning "using .num_workers \\(32\\)., ignoring 64" "" { target openacc_nvidia_accel_selected } l_compute$c_compute } */
{
int i, j;
int tmpvar;
int &tmpref = tmpvar;
#pragma acc loop collapse(2) gang worker vector private(tmpref) /* { dg-line l_loop[incr c_loop] } */
/* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
/* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
for (i = 0; i < 256; i++)
{
for (j = 0; j < 256; j++)
{
tmpref = (i * 256 + j) * 103;
res[i * 256 + j] = tmpref;
}
}
}
for (i = 0; i < 65536; i++)
if (res[i] != i * 103)
abort ();
}
int main (int argc, char *argv[])
{
gangs ();
workers ();
vectors ();
gangs_workers_vectors ();
return 0;
}
|
