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
|
// Ensure that a non-scalar dummy arguments which are implicitly used inside
// offloaded regions are properly mapped using present_or_copy semantics.
// { dg-do run }
#include <cassert>
const int n = 100;
struct data {
int v;
};
void
kernels_present (data &d, int &x)
{
#pragma acc kernels present (d, x) default (none)
{
d.v = x;
}
}
void
parallel_present (data &d, int &x)
{
#pragma acc parallel present (d, x) default (none)
{
d.v = x;
}
}
void
kernels_implicit (data &d, int &x)
{
#pragma acc kernels
{
d.v = x;
}
}
void
parallel_implicit (data &d, int &x)
{
#pragma acc parallel
{
d.v = x;
}
}
void
reference_data (data &d, int &x)
{
#pragma acc data copy(d, x)
{
kernels_present (d, x);
#pragma acc update host(d)
assert (d.v == x);
x = 200;
#pragma acc update device(x)
parallel_present (d, x);
}
assert (d.v == x);
x = 300;
kernels_implicit (d, x);
assert (d.v == x);
x = 400;
parallel_implicit (d, x);
assert (d.v == x);
}
int
main ()
{
data d;
int x = 100;
#pragma acc data copy(d, x)
{
kernels_present (d, x);
#pragma acc update host(d)
assert (d.v == x);
x = 200;
#pragma acc update device(x)
parallel_present (d, x);
}
assert (d.v == x);
x = 300;
kernels_implicit (d, x);
assert (d.v == x);
x = 400;
parallel_implicit (d, x);
assert (d.v == x);
reference_data (d, x);
return 0;
}
|
