blob: 2e48410b39dfb4e7c5eaa365074a125a4a823792 (
plain)
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
|
/* Exercise nested function decomposition, gcc/tree-nested.c. */
/* See gcc/testsuite/gfortran.dg/goacc/nested-function-1.f90 for the Fortran
version. */
/* { dg-additional-options "--param=openacc-kernels=decompose" } */
/* { dg-additional-options "-fopt-info-all-omp" } */
/* { dg-additional-options "--param=openacc-privatization=noisy" }
Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types):
{ dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
/* 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_loop 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". */
int main ()
{
#define N 100
int nonlocal_arg;
int nonlocal_a[N];
int nonlocal_i;
int nonlocal_j;
for (int i = 0; i < N; ++i)
nonlocal_a[i] = 5;
nonlocal_arg = 5;
void local ()
{
int local_i;
int local_arg;
int local_a[N];
int local_j;
for (int i = 0; i < N; ++i)
local_a[i] = 5;
local_arg = 5;
#pragma acc kernels loop /* { dg-line l_compute_loop[incr c_compute_loop] } */ \
gang(num:local_arg) worker(local_arg) vector(local_arg) \
wait async(local_arg)
/* { dg-note {OpenACC 'kernels' decomposition: variable 'local_arg' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop }
{ dg-note {variable 'local_arg' made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-note {variable 'local_arg\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-note {variable 'local_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-note {variable 'local_i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute_loop$c_compute_loop } */
for (local_i = 0; local_i < N; ++local_i)
{
#pragma acc cache (local_a[local_i:5])
local_a[local_i] = 100;
#pragma acc loop seq tile(*) /* { dg-line l_loop[incr c_loop] } */
/* { dg-note {variable 'local_j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop$c_loop } */
for (local_j = 0; local_j < N; ++local_j)
;
#pragma acc loop auto independent tile(1) /* { dg-line l_loop[incr c_loop] } */
/* { dg-note {variable 'local_j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop$c_loop } */
for (local_j = 0; local_j < N; ++local_j)
;
}
#pragma acc kernels loop /* { dg-line l_compute_loop[incr c_compute_loop] } */ \
gang(static:local_arg) worker(local_arg) vector(local_arg) \
wait(local_arg, local_arg + 1, local_arg + 2) async
/* { dg-note {OpenACC 'kernels' decomposition: variable 'local_arg' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop }
{ dg-note {variable 'local_arg' already made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-note {variable 'local_arg\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-note {variable 'local_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-note {variable 'local_i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute_loop$c_compute_loop } */
for (local_i = 0; local_i < N; ++local_i)
{
#pragma acc cache (local_a[local_i:4])
local_a[local_i] = 100;
#pragma acc loop seq tile(1) /* { dg-line l_loop[incr c_loop] } */
/* { dg-note {variable 'local_j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop$c_loop } */
for (local_j = 0; local_j < N; ++local_j)
;
#pragma acc loop auto independent tile(*) /* { dg-line l_loop[incr c_loop] } */
/* { dg-note {variable 'local_j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop$c_loop } */
for (local_j = 0; local_j < N; ++local_j)
;
}
}
void nonlocal ()
{
for (int i = 0; i < N; ++i)
nonlocal_a[i] = 5;
nonlocal_arg = 5;
#pragma acc kernels loop /* { dg-line l_compute_loop[incr c_compute_loop] } */ \
gang(num:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) \
wait async(nonlocal_arg)
/* { dg-note {OpenACC 'kernels' decomposition: variable 'nonlocal_arg' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop }
{ dg-note {variable 'nonlocal_arg' made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-note {variable 'nonlocal_arg\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-note {variable 'nonlocal_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-note {variable 'nonlocal_i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute_loop$c_compute_loop } */
for (nonlocal_i = 0; nonlocal_i < N; ++nonlocal_i)
{
#pragma acc cache (nonlocal_a[nonlocal_i:3])
nonlocal_a[nonlocal_i] = 100;
#pragma acc loop seq tile(2) /* { dg-line l_loop[incr c_loop] } */
/* { dg-note {variable 'nonlocal_j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop$c_loop } */
for (nonlocal_j = 0; nonlocal_j < N; ++nonlocal_j)
;
#pragma acc loop auto independent tile(3) /* { dg-line l_loop[incr c_loop] } */
/* { dg-note {variable 'nonlocal_j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop$c_loop } */
for (nonlocal_j = 0; nonlocal_j < N; ++nonlocal_j)
;
}
#pragma acc kernels loop /* { dg-line l_compute_loop[incr c_compute_loop] } */ \
gang(static:nonlocal_arg) worker(nonlocal_arg) vector(nonlocal_arg) \
wait(nonlocal_arg, nonlocal_arg + 1, nonlocal_arg + 2) async
/* { dg-note {OpenACC 'kernels' decomposition: variable 'nonlocal_arg' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop }
{ dg-note {variable 'nonlocal_arg' already made addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-note {variable 'nonlocal_arg\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-note {variable 'nonlocal_i\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-note {variable 'nonlocal_i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_compute_loop$c_compute_loop } */
/* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_compute_loop$c_compute_loop } */
for (nonlocal_i = 0; nonlocal_i < N; ++nonlocal_i)
{
#pragma acc cache (nonlocal_a[nonlocal_i:2])
nonlocal_a[nonlocal_i] = 100;
#pragma acc loop seq tile(*) /* { dg-line l_loop[incr c_loop] } */
/* { dg-note {variable 'nonlocal_j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop$c_loop } */
for (nonlocal_j = 0; nonlocal_j < N; ++nonlocal_j)
;
#pragma acc loop auto independent tile(*) /* { dg-line l_loop[incr c_loop] } */
/* { dg-note {variable 'nonlocal_j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop$c_loop } */
for (nonlocal_j = 0; nonlocal_j < N; ++nonlocal_j)
;
}
}
local ();
nonlocal ();
return 0;
}
|
