1 // RUN: %libomptarget-compile-generic -fopenmp-extensions
2 // RUN: %libomptarget-run-generic | %fcheck-generic -strict-whitespace
3 
4 #include <omp.h>
5 #include <stdio.h>
6 
7 #define CHECK_PRESENCE(Var1, Var2, Var3)                                       \
8   printf("    presence of %s, %s, %s: %d, %d, %d\n",                           \
9          #Var1, #Var2, #Var3,                                                  \
10          omp_target_is_present(&Var1, omp_get_default_device()),               \
11          omp_target_is_present(&Var2, omp_get_default_device()),               \
12          omp_target_is_present(&Var3, omp_get_default_device()))
13 
main()14 int main() {
15   int m, r, d;
16   // CHECK: presence of m, r, d: 0, 0, 0
17   CHECK_PRESENCE(m, r, d);
18 
19   // -----------------------------------------------------------------------
20   // CHECK-NEXT: check:{{.*}}
21   printf("check: dyn>0, hold=0, dec/reset dyn=0\n");
22 
23   // CHECK-NEXT: structured{{.*}}
24   printf("  structured dec of dyn\n");
25   #pragma omp target data map(tofrom: m) map(alloc: r, d)
26   {
27     // CHECK-NEXT: presence of m, r, d: 1, 1, 1
28     CHECK_PRESENCE(m, r, d);
29     #pragma omp target data map(tofrom: m) map(alloc: r, d)
30     {
31       // CHECK-NEXT: presence of m, r, d: 1, 1, 1
32       CHECK_PRESENCE(m, r, d);
33     }
34     // CHECK-NEXT: presence of m, r, d: 1, 1, 1
35     CHECK_PRESENCE(m, r, d);
36   }
37   // CHECK-NEXT: presence of m, r, d: 0, 0, 0
38   CHECK_PRESENCE(m, r, d);
39 
40   // CHECK-NEXT: dynamic{{.*}}
41   printf("  dynamic dec/reset of dyn\n");
42   #pragma omp target data map(tofrom: m) map(alloc: r, d)
43   {
44     // CHECK-NEXT: presence of m, r, d: 1, 1, 1
45     CHECK_PRESENCE(m, r, d);
46     #pragma omp target data map(tofrom: m) map(alloc: r, d)
47     {
48       // CHECK-NEXT: presence of m, r, d: 1, 1, 1
49       CHECK_PRESENCE(m, r, d);
50       #pragma omp target exit data map(from: m) map(release: r)
51       // CHECK-NEXT: presence of m, r, d: 1, 1, 1
52       CHECK_PRESENCE(m, r, d);
53       #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
54       // CHECK-NEXT: presence of m, r, d: 0, 0, 0
55       CHECK_PRESENCE(m, r, d);
56     }
57     // CHECK-NEXT: presence of m, r, d: 0, 0, 0
58     CHECK_PRESENCE(m, r, d);
59     #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
60     // CHECK-NEXT: presence of m, r, d: 0, 0, 0
61     CHECK_PRESENCE(m, r, d);
62   }
63   // CHECK-NEXT: presence of m, r, d: 0, 0, 0
64   CHECK_PRESENCE(m, r, d);
65 
66   // -----------------------------------------------------------------------
67   // CHECK: check:{{.*}}
68   printf("check: dyn=0, hold>0, dec/reset dyn=0, dec hold=0\n");
69 
70   // Structured dec of dyn would require dyn>0.
71 
72   // CHECK-NEXT: dynamic{{.*}}
73   printf("  dynamic dec/reset of dyn\n");
74   #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
75   {
76     // CHECK-NEXT: presence of m, r, d: 1, 1, 1
77     CHECK_PRESENCE(m, r, d);
78     #pragma omp target data map(ompx_hold, tofrom: m) \
79                             map(ompx_hold, alloc: r, d)
80     {
81       // CHECK-NEXT: presence of m, r, d: 1, 1, 1
82       CHECK_PRESENCE(m, r, d);
83       #pragma omp target exit data map(from: m) map(release: r)
84       // CHECK-NEXT: presence of m, r, d: 1, 1, 1
85       CHECK_PRESENCE(m, r, d);
86       #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
87       // CHECK-NEXT: presence of m, r, d: 1, 1, 1
88       CHECK_PRESENCE(m, r, d);
89     }
90     // CHECK-NEXT: presence of m, r, d: 1, 1, 1
91     CHECK_PRESENCE(m, r, d);
92     #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
93     // CHECK-NEXT: presence of m, r, d: 1, 1, 1
94     CHECK_PRESENCE(m, r, d);
95   }
96   // CHECK-NEXT: presence of m, r, d: 0, 0, 0
97   CHECK_PRESENCE(m, r, d);
98 
99   // -----------------------------------------------------------------------
100   // CHECK: check:{{.*}}
101   printf("check: dyn>0, hold>0, dec/reset dyn=0, dec hold=0\n");
102 
103   // CHECK-NEXT: structured{{.*}}
104   printf("  structured dec of dyn\n");
105   #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
106   {
107     // CHECK-NEXT: presence of m, r, d: 1, 1, 1
108     CHECK_PRESENCE(m, r, d);
109     #pragma omp target data map(ompx_hold, tofrom: m) \
110                             map(ompx_hold, alloc: r, d)
111     {
112       // CHECK-NEXT: presence of m, r, d: 1, 1, 1
113       CHECK_PRESENCE(m, r, d);
114       #pragma omp target data map(tofrom: m) map(alloc: r, d)
115       {
116         // CHECK-NEXT: presence of m, r, d: 1, 1, 1
117         CHECK_PRESENCE(m, r, d);
118         #pragma omp target data map(tofrom: m) map(alloc: r, d)
119         {
120           // CHECK-NEXT: presence of m, r, d: 1, 1, 1
121           CHECK_PRESENCE(m, r, d);
122         }
123         // CHECK-NEXT: presence of m, r, d: 1, 1, 1
124         CHECK_PRESENCE(m, r, d);
125       }
126       // CHECK-NEXT: presence of m, r, d: 1, 1, 1
127       CHECK_PRESENCE(m, r, d);
128     }
129     // CHECK-NEXT: presence of m, r, d: 1, 1, 1
130     CHECK_PRESENCE(m, r, d);
131   }
132   // CHECK-NEXT: presence of m, r, d: 0, 0, 0
133   CHECK_PRESENCE(m, r, d);
134 
135   // CHECK-NEXT: dynamic{{.*}}
136   printf("  dynamic dec/reset of dyn\n");
137   #pragma omp target enter data map(to: m) map(alloc: r, d)
138   // CHECK-NEXT: presence of m, r, d: 1, 1, 1
139   CHECK_PRESENCE(m, r, d);
140   #pragma omp target enter data map(to: m) map(alloc: r, d)
141   // CHECK-NEXT: presence of m, r, d: 1, 1, 1
142   CHECK_PRESENCE(m, r, d);
143   #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
144   {
145     // CHECK-NEXT: presence of m, r, d: 1, 1, 1
146     CHECK_PRESENCE(m, r, d);
147     #pragma omp target data map(ompx_hold, tofrom: m) \
148                             map(ompx_hold, alloc: r, d)
149     {
150       // CHECK-NEXT: presence of m, r, d: 1, 1, 1
151       CHECK_PRESENCE(m, r, d);
152       #pragma omp target exit data map(from: m) map(release: r)
153       // CHECK-NEXT: presence of m, r, d: 1, 1, 1
154       CHECK_PRESENCE(m, r, d);
155       #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
156       // CHECK-NEXT: presence of m, r, d: 1, 1, 1
157       CHECK_PRESENCE(m, r, d);
158     }
159     // CHECK-NEXT: presence of m, r, d: 1, 1, 1
160     CHECK_PRESENCE(m, r, d);
161     #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
162     // CHECK-NEXT: presence of m, r, d: 1, 1, 1
163     CHECK_PRESENCE(m, r, d);
164   }
165   // CHECK-NEXT: presence of m, r, d: 0, 0, 0
166   CHECK_PRESENCE(m, r, d);
167 
168   // -----------------------------------------------------------------------
169   // CHECK: check:{{.*}}
170   printf("check: dyn>0, hold>0, dec hold=0, dec/reset dyn=0\n");
171 
172   // CHECK-NEXT: structured{{.*}}
173   printf("  structured dec of dyn\n");
174   #pragma omp target data map(tofrom: m) map(alloc: r, d)
175   {
176     // CHECK-NEXT: presence of m, r, d: 1, 1, 1
177     CHECK_PRESENCE(m, r, d);
178     #pragma omp target data map(tofrom: m) map(alloc: r, d)
179     {
180       // CHECK-NEXT: presence of m, r, d: 1, 1, 1
181       CHECK_PRESENCE(m, r, d);
182       #pragma omp target data map(ompx_hold, tofrom: m) \
183                               map(ompx_hold, alloc: r, d)
184       {
185         // CHECK-NEXT: presence of m, r, d: 1, 1, 1
186         CHECK_PRESENCE(m, r, d);
187         #pragma omp target data map(ompx_hold, tofrom: m) \
188                                 map(ompx_hold, alloc: r, d)
189         {
190           // CHECK-NEXT: presence of m, r, d: 1, 1, 1
191           CHECK_PRESENCE(m, r, d);
192         }
193         // CHECK-NEXT: presence of m, r, d: 1, 1, 1
194         CHECK_PRESENCE(m, r, d);
195       }
196       // CHECK-NEXT: presence of m, r, d: 1, 1, 1
197       CHECK_PRESENCE(m, r, d);
198     }
199     // CHECK-NEXT: presence of m, r, d: 1, 1, 1
200     CHECK_PRESENCE(m, r, d);
201   }
202   // CHECK-NEXT: presence of m, r, d: 0, 0, 0
203   CHECK_PRESENCE(m, r, d);
204 
205   // CHECK-NEXT: dynamic{{.*}}
206   printf("  dynamic dec/reset of dyn\n");
207   #pragma omp target enter data map(to: m) map(alloc: r, d)
208   // CHECK-NEXT: presence of m, r, d: 1, 1, 1
209   CHECK_PRESENCE(m, r, d);
210   #pragma omp target enter data map(to: m) map(alloc: r, d)
211   // CHECK-NEXT: presence of m, r, d: 1, 1, 1
212   CHECK_PRESENCE(m, r, d);
213   #pragma omp target data map(ompx_hold, tofrom: m) map(ompx_hold, alloc: r, d)
214   {
215     // CHECK-NEXT: presence of m, r, d: 1, 1, 1
216     CHECK_PRESENCE(m, r, d);
217     #pragma omp target data map(ompx_hold, tofrom: m) \
218                             map(ompx_hold, alloc: r, d)
219     {
220       // CHECK-NEXT: presence of m, r, d: 1, 1, 1
221       CHECK_PRESENCE(m, r, d);
222     }
223     // CHECK-NEXT: presence of m, r, d: 1, 1, 1
224     CHECK_PRESENCE(m, r, d);
225   }
226   // CHECK-NEXT: presence of m, r, d: 1, 1, 1
227   CHECK_PRESENCE(m, r, d);
228   #pragma omp target exit data map(from: m) map(release: r)
229   // CHECK-NEXT: presence of m, r, d: 1, 1, 1
230   CHECK_PRESENCE(m, r, d);
231   #pragma omp target exit data map(from: m) map(release: r) map(delete: d)
232   // CHECK-NEXT: presence of m, r, d: 0, 0, 0
233   CHECK_PRESENCE(m, r, d);
234 
235   return 0;
236 }
237