1 /* Test transitioning of data lifetimes between structured and dynamic. */
2
3 /* { dg-skip-if "" { *-*-* } { "-DACC_MEM_SHARED=1" } } */
4
5 #include <openacc.h>
6 #include <assert.h>
7 #include <stdlib.h>
8
9 #define SIZE 1024
10
11 void
f1(void)12 f1 (void)
13 {
14 char *block1 = (char *) malloc (SIZE);
15
16 #ifdef OPENACC_API
17 acc_copyin (block1, SIZE);
18 acc_copyin (block1, SIZE);
19 #else
20 #pragma acc enter data copyin(block1[0:SIZE])
21 #pragma acc enter data copyin(block1[0:SIZE])
22 #endif
23
24 #pragma acc data copy(block1[0:SIZE])
25 {
26 #ifdef OPENACC_API
27 acc_copyin (block1, SIZE);
28 #else
29 #pragma acc enter data copyin(block1[0:SIZE])
30 #endif
31 }
32
33 assert (acc_is_present (block1, SIZE));
34
35 #ifdef OPENACC_API
36 acc_copyout (block1, SIZE);
37 assert (acc_is_present (block1, SIZE));
38 acc_copyout (block1, SIZE);
39 assert (acc_is_present (block1, SIZE));
40 acc_copyout (block1, SIZE);
41 assert (!acc_is_present (block1, SIZE));
42 #else
43 #pragma acc exit data copyout(block1[0:SIZE])
44 assert (acc_is_present (block1, SIZE));
45 #pragma acc exit data copyout(block1[0:SIZE])
46 assert (acc_is_present (block1, SIZE));
47 #pragma acc exit data copyout(block1[0:SIZE])
48 assert (!acc_is_present (block1, SIZE));
49 #endif
50
51 free (block1);
52 }
53
54 void
f2(void)55 f2 (void)
56 {
57 char *block1 = (char *) malloc (SIZE);
58
59 #ifdef OPENACC_API
60 acc_copyin (block1, SIZE);
61 #else
62 #pragma acc enter data copyin(block1[0:SIZE])
63 #endif
64
65 #pragma acc data copy(block1[0:SIZE])
66 {
67 #ifdef OPENACC_API
68 acc_copyout (block1, SIZE);
69 #else
70 #pragma acc exit data copyout(block1[0:SIZE])
71 #endif
72 /* This should stay present until the end of the structured data
73 lifetime. */
74 assert (acc_is_present (block1, SIZE));
75 }
76
77 assert (!acc_is_present (block1, SIZE));
78
79 free (block1);
80 }
81
82 void
f3(void)83 f3 (void)
84 {
85 char *block1 = (char *) malloc (SIZE);
86
87 #ifdef OPENACC_API
88 acc_copyin (block1, SIZE);
89 #else
90 #pragma acc enter data copyin(block1[0:SIZE])
91 #endif
92
93 #pragma acc data copy(block1[0:SIZE])
94 {
95 #ifdef OPENACC_API
96 acc_copyout (block1, SIZE);
97 acc_copyin (block1, SIZE);
98 #else
99 #pragma acc exit data copyout(block1[0:SIZE])
100 #pragma acc enter data copyin(block1[0:SIZE])
101 #endif
102 assert (acc_is_present (block1, SIZE));
103 }
104
105 assert (acc_is_present (block1, SIZE));
106 #ifdef OPENACC_API
107 acc_copyout (block1, SIZE);
108 #else
109 #pragma acc exit data copyout(block1[0:SIZE])
110 #endif
111 assert (!acc_is_present (block1, SIZE));
112
113 free (block1);
114 }
115
116 void
f4(void)117 f4 (void)
118 {
119 char *block1 = (char *) malloc (SIZE);
120 char *block2 = (char *) malloc (SIZE);
121 char *block3 = (char *) malloc (SIZE);
122
123 #pragma acc data copy(block1[0:SIZE], block2[0:SIZE], block3[0:SIZE])
124 {
125 /* The first copyin of block2 is the enclosing data region. This
126 "enter data" should make it live beyond the end of this region.
127 This works, though the on-target copies of block1, block2 and block3
128 will stay allocated until block2 is unmapped because they are bound
129 together in a single target_mem_desc. */
130 #ifdef OPENACC_API
131 acc_copyin (block2, SIZE);
132 #else
133 #pragma acc enter data copyin(block2[0:SIZE])
134 #endif
135 }
136
137 assert (!acc_is_present (block1, SIZE));
138 assert (acc_is_present (block2, SIZE));
139 assert (!acc_is_present (block3, SIZE));
140
141 #ifdef OPENACC_API
142 acc_copyout (block2, SIZE);
143 #else
144 #pragma acc exit data copyout(block2[0:SIZE])
145 #endif
146 assert (!acc_is_present (block2, SIZE));
147
148 free (block1);
149 free (block2);
150 free (block3);
151 }
152
153 int
main(int argc,char * argv[])154 main (int argc, char *argv[])
155 {
156 f1 ();
157 f2 ();
158 f3 ();
159 f4 ();
160 return 0;
161 }
162