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
|
/* Test transitioning of data lifetimes between structured and dynamic. */
/* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
#include <openacc.h>
#include <assert.h>
#include <stdlib.h>
#define SIZE 1024
void
f1 (void)
{
char *block1 = (char *) malloc (SIZE);
#ifdef OPENACC_API
acc_copyin (block1, SIZE);
acc_copyin (block1, SIZE);
#else
#pragma acc enter data copyin(block1[0:SIZE])
#pragma acc enter data copyin(block1[0:SIZE])
#endif
#pragma acc data copy(block1[0:SIZE])
{
#ifdef OPENACC_API
acc_copyin (block1, SIZE);
#else
#pragma acc enter data copyin(block1[0:SIZE])
#endif
}
assert (acc_is_present (block1, SIZE));
#ifdef OPENACC_API
acc_copyout (block1, SIZE);
assert (acc_is_present (block1, SIZE));
acc_copyout (block1, SIZE);
assert (acc_is_present (block1, SIZE));
acc_copyout (block1, SIZE);
assert (!acc_is_present (block1, SIZE));
#else
#pragma acc exit data copyout(block1[0:SIZE])
assert (acc_is_present (block1, SIZE));
#pragma acc exit data copyout(block1[0:SIZE])
assert (acc_is_present (block1, SIZE));
#pragma acc exit data copyout(block1[0:SIZE])
assert (!acc_is_present (block1, SIZE));
#endif
free (block1);
}
void
f2 (void)
{
char *block1 = (char *) malloc (SIZE);
#ifdef OPENACC_API
acc_copyin (block1, SIZE);
#else
#pragma acc enter data copyin(block1[0:SIZE])
#endif
#pragma acc data copy(block1[0:SIZE])
{
#ifdef OPENACC_API
acc_copyout (block1, SIZE);
#else
#pragma acc exit data copyout(block1[0:SIZE])
#endif
/* This should stay present until the end of the structured data
lifetime. */
assert (acc_is_present (block1, SIZE));
}
assert (!acc_is_present (block1, SIZE));
free (block1);
}
void
f3 (void)
{
char *block1 = (char *) malloc (SIZE);
#ifdef OPENACC_API
acc_copyin (block1, SIZE);
#else
#pragma acc enter data copyin(block1[0:SIZE])
#endif
#pragma acc data copy(block1[0:SIZE])
{
#ifdef OPENACC_API
acc_copyout (block1, SIZE);
acc_copyin (block1, SIZE);
#else
#pragma acc exit data copyout(block1[0:SIZE])
#pragma acc enter data copyin(block1[0:SIZE])
#endif
assert (acc_is_present (block1, SIZE));
}
assert (acc_is_present (block1, SIZE));
#ifdef OPENACC_API
acc_copyout (block1, SIZE);
#else
#pragma acc exit data copyout(block1[0:SIZE])
#endif
assert (!acc_is_present (block1, SIZE));
free (block1);
}
void
f4 (void)
{
char *block1 = (char *) malloc (SIZE);
char *block2 = (char *) malloc (SIZE);
char *block3 = (char *) malloc (SIZE);
#pragma acc data copy(block1[0:SIZE], block2[0:SIZE], block3[0:SIZE])
{
/* The first copyin of block2 is the enclosing data region. This
"enter data" should make it live beyond the end of this region.
This works, though the on-target copies of block1, block2 and block3
will stay allocated until block2 is unmapped because they are bound
together in a single target_mem_desc. */
#ifdef OPENACC_API
acc_copyin (block2, SIZE);
#else
#pragma acc enter data copyin(block2[0:SIZE])
#endif
}
assert (!acc_is_present (block1, SIZE));
assert (acc_is_present (block2, SIZE));
assert (!acc_is_present (block3, SIZE));
#ifdef OPENACC_API
acc_copyout (block2, SIZE);
#else
#pragma acc exit data copyout(block2[0:SIZE])
#endif
assert (!acc_is_present (block2, SIZE));
free (block1);
free (block2);
free (block3);
}
int
main (int argc, char *argv[])
{
f1 ();
f2 ();
f3 ();
f4 ();
return 0;
}
|
